[clang] [llvm] [Clang][OpenMP] Support expression semantics in target update fields with non-contiguous array sections (PR #176708)
Amit Tiwari via llvm-commits
llvm-commits at lists.llvm.org
Mon Feb 16 10:29:45 PST 2026
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/176708
>From ee31e77170f8a6970f9360a352d25ba97083bbc5 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/6] 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 464ec5b5a2ece..44435128b3242 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9825,16 +9825,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 8258eba748ee81a544e4ceb689b18af950c28583 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/6] 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 3f3c239c2de78..d983a28bc11d1 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8056,12 +8056,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 119439d82702fa733539d38056c31f540efe4720 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/6] 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 144ee997a2ea5ce3e792109d731a95a67eed26b5 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/6] 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 d983a28bc11d1..df4498d763b6e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8056,9 +8056,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 44435128b3242..067af7405cdad 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9828,8 +9828,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
>From 2b22e7a6c1cf9cbaa0ada34e12df5e13dd7b01aa Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 16 Feb 2026 10:14:28 -0500
Subject: [PATCH 5/6] comments_addressed
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 10 +-
...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 -----
.../offloading/strided_multiple_update_from.c | 11 +-
.../offloading/strided_partial_update_from.c | 8 +-
.../strided_ptr_multiple_update_from.c | 11 +-
.../strided_ptr_partial_update_from.c | 8 +-
.../strided_update_count_expression.c | 25 +-
.../strided_update_count_expression_complex.c | 46 ++-
.../strided_update_count_expression_misc.c | 32 +-
offload/test/offloading/strided_update_from.c | 8 +-
..._update_multiple_arrays_count_expression.c | 17 +-
...d_update_multiple_arrays_variable_stride.c | 28 +-
...strided_update_variable_count_and_stride.c | 26 +-
.../strided_update_variable_stride.c | 25 +-
.../strided_update_variable_stride_complex.c | 293 ------------------
.../strided_update_variable_stride_misc.c | 32 +-
20 files changed, 128 insertions(+), 792 deletions(-)
delete mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c
delete mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c
delete mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c
delete mode 100644 clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c
delete mode 100644 clang/test/OpenMP/target_update_variable_count_and_stride_messages.c
delete mode 100644 offload/test/offloading/strided_update_variable_stride_complex.c
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index df4498d763b6e..0098b9c64ae3e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8060,13 +8060,9 @@ class MappableExprsHandler {
assert(StrideExpr->getType()->isIntegerType() &&
"Stride expression must be of integer type");
- // If the stride is a variable (not a constant), it's non-contiguous.
+ // If the stride involves member access or array subscript, 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;
@@ -8075,7 +8071,7 @@ class MappableExprsHandler {
const auto Constant =
StrideExpr->getIntegerConstantExpr(CGF.getContext());
if (!Constant)
- return false;
+ return true;
// Treat non-unitary strides as non-contiguous.
return !Constant->isOne();
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
deleted file mode 100644
index 932cd6b1c97bb..0000000000000
--- a/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c
+++ /dev/null
@@ -1,62 +0,0 @@
-// 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
deleted file mode 100644
index 23fba9c8bc84f..0000000000000
--- a/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c
+++ /dev/null
@@ -1,57 +0,0 @@
-// 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
deleted file mode 100644
index 3f85ed0c48d66..0000000000000
--- a/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c
+++ /dev/null
@@ -1,64 +0,0 @@
-// 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
deleted file mode 100644
index 70775d5c8322c..0000000000000
--- a/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c
+++ /dev/null
@@ -1,72 +0,0 @@
-// 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
deleted file mode 100644
index 0082539538a32..0000000000000
--- a/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c
+++ /dev/null
@@ -1,85 +0,0 @@
-// 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_multiple_update_from.c b/offload/test/offloading/strided_multiple_update_from.c
index 4f2df81933738..7faac2f416260 100644
--- a/offload/test/offloading/strided_multiple_update_from.c
+++ b/offload/test/offloading/strided_multiple_update_from.c
@@ -11,13 +11,10 @@ int main() {
int len = 12;
double data1[len], data2[len];
-// Initial values
-#pragma omp target map(tofrom : data1[0 : len], data2[0 : len])
- {
- for (int i = 0; i < len; i++) {
- data1[i] = i;
- data2[i] = i * 10;
- }
+ // Initialize data on host
+ for (int i = 0; i < len; i++) {
+ data1[i] = i;
+ data2[i] = i * 10;
}
printf("original host array values:\n");
diff --git a/offload/test/offloading/strided_partial_update_from.c b/offload/test/offloading/strided_partial_update_from.c
index 4a2977a08f70b..7e491dcc0a883 100644
--- a/offload/test/offloading/strided_partial_update_from.c
+++ b/offload/test/offloading/strided_partial_update_from.c
@@ -11,11 +11,9 @@ int main() {
int len = 11;
double data[len];
-#pragma omp target map(tofrom : data[0 : len])
- {
- for (int i = 0; i < len; i++)
- data[i] = i;
- }
+ // Initialize data on host
+ for (int i = 0; i < len; i++)
+ data[i] = i;
// Initial values
printf("original host array values:\n");
diff --git a/offload/test/offloading/strided_ptr_multiple_update_from.c b/offload/test/offloading/strided_ptr_multiple_update_from.c
index 33deb59534f69..0bb3e62c48a1b 100644
--- a/offload/test/offloading/strided_ptr_multiple_update_from.c
+++ b/offload/test/offloading/strided_ptr_multiple_update_from.c
@@ -13,13 +13,10 @@ int main() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < len; i++) {
+ data1[i] = i;
+ data2[i] = i * 10;
}
printf("original host array values:\n");
diff --git a/offload/test/offloading/strided_ptr_partial_update_from.c b/offload/test/offloading/strided_ptr_partial_update_from.c
index be235d3d9c532..8d45ee05da5e0 100644
--- a/offload/test/offloading/strided_ptr_partial_update_from.c
+++ b/offload/test/offloading/strided_ptr_partial_update_from.c
@@ -12,11 +12,9 @@ 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;
- }
+ // Initialize data on host
+ for (int i = 0; i < len; i++)
+ data[i] = i;
// Initial values
printf("original host array values:\n");
diff --git a/offload/test/offloading/strided_update_count_expression.c b/offload/test/offloading/strided_update_count_expression.c
index a87da289a9154..e18fc881a8021 100644
--- a/offload/test/offloading/strided_update_count_expression.c
+++ b/offload/test/offloading/strided_update_count_expression.c
@@ -7,19 +7,13 @@
#include <omp.h>
#include <stdio.h>
-int main() {
+void test_1_update_from() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
}
printf("Test 1: Update FROM device\n");
@@ -40,10 +34,11 @@ int main() {
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)
- // ====================================================================
+void test_2_update_to() {
+ int len = 10;
+ double data[len];
for (int i = 0; i < len; i++) {
data[i] = i;
@@ -78,7 +73,11 @@ int main() {
printf("device array values after update to:\n");
for (int i = 0; i < len; i++)
printf("%f\n", data[i]);
+}
+int main() {
+ test_1_update_from();
+ test_2_update_to();
return 0;
}
diff --git a/offload/test/offloading/strided_update_count_expression_complex.c b/offload/test/offloading/strided_update_count_expression_complex.c
index f9beef513da24..560c4e767f882 100644
--- a/offload/test/offloading/strided_update_count_expression_complex.c
+++ b/offload/test/offloading/strided_update_count_expression_complex.c
@@ -11,10 +11,8 @@ struct Data {
double arr[20];
};
-int main() {
+void test_1_complex_count_expressions() {
struct Data s1, s2;
-
- // Test 1: Multiple arrays with different count expressions
s1.len = 10;
s2.len = 10;
@@ -56,15 +54,12 @@ int main() {
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;
- }
+ // Reset for TO test - initialize on host
+ 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
@@ -100,8 +95,10 @@ int main() {
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
+void test_2_complex_count_with_offset() {
+ struct Data s1, s2;
s1.offset = 2;
s1.len = 10;
s2.offset = 1;
@@ -146,15 +143,12 @@ int main() {
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;
- }
+ // Reset for TO test - initialize on host
+ 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
@@ -191,8 +185,6 @@ int main() {
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):
@@ -287,3 +279,9 @@ int main() {
// CHECK-NEXT: 53.000000
// CHECK-NEXT: 160.000000
// CHECK-NEXT: 180.000000
+
+int main() {
+ test_1_complex_count_expressions();
+ test_2_complex_count_with_offset();
+ return 0;
+}
diff --git a/offload/test/offloading/strided_update_count_expression_misc.c b/offload/test/offloading/strided_update_count_expression_misc.c
index 0e93a6d7df2cb..6c13eae4d37e2 100644
--- a/offload/test/offloading/strided_update_count_expression_misc.c
+++ b/offload/test/offloading/strided_update_count_expression_misc.c
@@ -6,20 +6,14 @@
#include <omp.h>
#include <stdio.h>
-int main() {
- // ====================================================================
- // TEST 1: Modulo operation in count expression
- // ====================================================================
-
+void test_1_modulo_count() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < len1; i++) {
+ data1[i] = i;
}
#pragma omp target data map(to : len1, divisor, data1[0 : len1])
@@ -38,20 +32,16 @@ int main() {
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
- // ====================================================================
-
+void test_2_large_stride_count() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < len2; i++) {
+ data2[i] = i;
}
#pragma omp target data map(to : len2, stride, data2[0 : len2])
@@ -70,7 +60,11 @@ int main() {
printf("\nTest 2: Large stride count expression\n");
for (int i = 0; i < len2; i++)
printf("%f\n", data2[i]);
+}
+int main() {
+ test_1_modulo_count();
+ test_2_large_stride_count();
return 0;
}
diff --git a/offload/test/offloading/strided_update_from.c b/offload/test/offloading/strided_update_from.c
index 9910bed9c5466..a019a025355a6 100644
--- a/offload/test/offloading/strided_update_from.c
+++ b/offload/test/offloading/strided_update_from.c
@@ -11,11 +11,9 @@
int main() {
int len = 8;
double data[len];
-#pragma omp target map(tofrom : len, data[0 : len])
- {
- for (int i = 0; i < len; i++) {
- data[i] = i;
- }
+ // Initialize data on host
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
}
// Initial values
printf("original host array values:\n");
diff --git a/offload/test/offloading/strided_update_multiple_arrays_count_expression.c b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c
index 9449baa663f67..a1472cacc4a39 100644
--- a/offload/test/offloading/strided_update_multiple_arrays_count_expression.c
+++ b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c
@@ -7,14 +7,10 @@
#include <omp.h>
#include <stdio.h>
-int main() {
+void test_1_update_from_multiple() {
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++) {
@@ -52,10 +48,11 @@ int main() {
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
- // ====================================================================
+void test_2_update_to_multiple() {
+ int n1 = 10, n2 = 10;
+ double arr1[n1], arr2[n2];
for (int i = 0; i < n1; i++) {
arr1[i] = i;
@@ -106,7 +103,11 @@ int main() {
printf("\ndevice arr2 values after update to:\n");
for (int i = 0; i < n2; i++)
printf("%f\n", arr2[i]);
+}
+int main() {
+ test_1_update_from_multiple();
+ test_2_update_to_multiple();
return 0;
}
diff --git a/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c
index 68c3eca4ccc56..32859202a20fa 100644
--- a/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c
+++ b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c
@@ -5,21 +5,15 @@
#include <omp.h>
#include <stdio.h>
-int main() {
+void test_1_update_from_multiple() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < 10; i++) {
+ data1[i] = i;
+ data2[i] = i * 10;
}
printf("Test 1: Update FROM - Multiple arrays\n");
@@ -44,10 +38,12 @@ int main() {
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
- // ====================================================================
+void test_2_update_to_multiple() {
+ int stride1 = 2;
+ int stride2 = 2;
+ double data1[10], data2[10];
for (int i = 0; i < 10; i++) {
data1[i] = i;
@@ -90,7 +86,11 @@ int main() {
printf("\ndevice data2 after update to:\n");
for (int i = 0; i < 10; i++)
printf("%f\n", data2[i]);
+}
+int main() {
+ test_1_update_from_multiple();
+ test_2_update_to_multiple();
return 0;
}
diff --git a/offload/test/offloading/strided_update_variable_count_and_stride.c b/offload/test/offloading/strided_update_variable_count_and_stride.c
index 36056ab64250a..1e1e41653c2c4 100644
--- a/offload/test/offloading/strided_update_variable_count_and_stride.c
+++ b/offload/test/offloading/strided_update_variable_count_and_stride.c
@@ -5,20 +5,14 @@
#include <omp.h>
#include <stdio.h>
-int main() {
+void test_1_update_from() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
}
printf("Test 1: Update FROM - Variable count and stride\n");
@@ -41,10 +35,12 @@ int main() {
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
- // ====================================================================
+void test_2_update_to() {
+ int len = 10;
+ int stride = 2;
+ double data[len];
for (int i = 0; i < len; i++) {
data[i] = i;
@@ -81,7 +77,11 @@ int main() {
printf("device values after update to:\n");
for (int i = 0; i < len; i++)
printf("%f\n", data[i]);
+}
+int main() {
+ test_1_update_from();
+ test_2_update_to();
return 0;
}
diff --git a/offload/test/offloading/strided_update_variable_stride.c b/offload/test/offloading/strided_update_variable_stride.c
index 94723d91734a6..7c8079efa56ea 100644
--- a/offload/test/offloading/strided_update_variable_stride.c
+++ b/offload/test/offloading/strided_update_variable_stride.c
@@ -5,19 +5,13 @@
#include <omp.h>
#include <stdio.h>
-int main() {
+void test_1_update_from() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < 10; i++) {
+ data[i] = i;
}
printf("Test 1: Update FROM device\n");
@@ -40,10 +34,11 @@ int main() {
printf("from target results:\n");
for (int i = 0; i < 10; i++)
printf("%f\n", data[i]);
+}
- // ====================================================================
- // TEST 2: Update TO device (host -> device)
- // ====================================================================
+void test_2_update_to() {
+ int stride = 2;
+ double data[10];
for (int i = 0; i < 10; i++) {
data[i] = i;
@@ -80,7 +75,11 @@ int main() {
printf("device values after update to:\n");
for (int i = 0; i < 10; i++)
printf("%f\n", data[i]);
+}
+int main() {
+ test_1_update_from();
+ test_2_update_to();
return 0;
}
diff --git a/offload/test/offloading/strided_update_variable_stride_complex.c b/offload/test/offloading/strided_update_variable_stride_complex.c
deleted file mode 100644
index 3c9857ec22178..0000000000000
--- a/offload/test/offloading/strided_update_variable_stride_complex.c
+++ /dev/null
@@ -1,293 +0,0 @@
-// 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
index d27ae0123bfa8..df2bfa64cfe3c 100644
--- a/offload/test/offloading/strided_update_variable_stride_misc.c
+++ b/offload/test/offloading/strided_update_variable_stride_misc.c
@@ -5,19 +5,13 @@
#include <omp.h>
#include <stdio.h>
-int main() {
- // ====================================================================
- // TEST 1: Variable stride = 1 (contiguous, but detected as variable)
- // ====================================================================
-
+void test_1_variable_stride_one() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < 10; i++) {
+ data1[i] = i;
}
#pragma omp target data map(to : stride_one, data1[0 : 10])
@@ -35,19 +29,15 @@ int main() {
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)
- // ====================================================================
-
+void test_2_variable_stride_large() {
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;
- }
+ // Initialize data on host
+ for (int i = 0; i < 10; i++) {
+ data2[i] = i;
}
#pragma omp target data map(to : stride_large, data2[0 : 10])
@@ -65,7 +55,11 @@ int main() {
printf("\nTest 2: Variable stride = 5\n");
for (int i = 0; i < 10; i++)
printf("%f\n", data2[i]);
+}
+int main() {
+ test_1_variable_stride_one();
+ test_2_variable_stride_large();
return 0;
}
>From 5fff77e56bbfe9e986f1125e3503abf2f738dddd Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 16 Feb 2026 13:26:41 -0500
Subject: [PATCH 6/6] cleanup
---
.../test/offloading/strided_multiple_update_from.c | 11 +++++++----
offload/test/offloading/strided_partial_update_from.c | 8 +++++---
.../offloading/strided_ptr_multiple_update_from.c | 11 +++++++----
.../test/offloading/strided_ptr_partial_update_from.c | 8 +++++---
offload/test/offloading/strided_update_from.c | 8 +++++---
5 files changed, 29 insertions(+), 17 deletions(-)
diff --git a/offload/test/offloading/strided_multiple_update_from.c b/offload/test/offloading/strided_multiple_update_from.c
index 7faac2f416260..4f2df81933738 100644
--- a/offload/test/offloading/strided_multiple_update_from.c
+++ b/offload/test/offloading/strided_multiple_update_from.c
@@ -11,10 +11,13 @@ int main() {
int len = 12;
double data1[len], data2[len];
- // Initialize data on host
- for (int i = 0; i < len; i++) {
- data1[i] = i;
- data2[i] = i * 10;
+// 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");
diff --git a/offload/test/offloading/strided_partial_update_from.c b/offload/test/offloading/strided_partial_update_from.c
index 7e491dcc0a883..4a2977a08f70b 100644
--- a/offload/test/offloading/strided_partial_update_from.c
+++ b/offload/test/offloading/strided_partial_update_from.c
@@ -11,9 +11,11 @@ int main() {
int len = 11;
double data[len];
- // Initialize data on host
- for (int i = 0; i < len; i++)
- data[i] = i;
+#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");
diff --git a/offload/test/offloading/strided_ptr_multiple_update_from.c b/offload/test/offloading/strided_ptr_multiple_update_from.c
index 0bb3e62c48a1b..33deb59534f69 100644
--- a/offload/test/offloading/strided_ptr_multiple_update_from.c
+++ b/offload/test/offloading/strided_ptr_multiple_update_from.c
@@ -13,10 +13,13 @@ int main() {
double *data1 = (double *)calloc(len, sizeof(double));
double *data2 = (double *)calloc(len, sizeof(double));
- // Initialize data on host
- for (int i = 0; i < len; i++) {
- data1[i] = i;
- data2[i] = i * 10;
+// 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");
diff --git a/offload/test/offloading/strided_ptr_partial_update_from.c b/offload/test/offloading/strided_ptr_partial_update_from.c
index 8d45ee05da5e0..be235d3d9c532 100644
--- a/offload/test/offloading/strided_ptr_partial_update_from.c
+++ b/offload/test/offloading/strided_ptr_partial_update_from.c
@@ -12,9 +12,11 @@ int main() {
int len = 11;
double *data = (double *)calloc(len, sizeof(double));
- // Initialize data on host
- for (int i = 0; i < len; i++)
- data[i] = i;
+#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");
diff --git a/offload/test/offloading/strided_update_from.c b/offload/test/offloading/strided_update_from.c
index a019a025355a6..9910bed9c5466 100644
--- a/offload/test/offloading/strided_update_from.c
+++ b/offload/test/offloading/strided_update_from.c
@@ -11,9 +11,11 @@
int main() {
int len = 8;
double data[len];
- // Initialize data on host
- for (int i = 0; i < len; i++) {
- data[i] = i;
+#pragma omp target map(tofrom : len, data[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ }
}
// Initial values
printf("original host array values:\n");
More information about the llvm-commits
mailing list