[clang] [llvm] [Clang][OpenMP] Support expression semantics in target update fields with non-contiguous array sections (PR #176708)
Amit Tiwari via cfe-commits
cfe-commits at lists.llvm.org
Mon Mar 2 00:19:18 PST 2026
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/176708
>From 5539922bb5f35df55fa8f2fcf0f1160427c6ab49 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/9] 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 25f4da7c90d90..da03864cd6f15 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9945,16 +9945,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 f9643a6d602ddc4ac9e2b24c31f226ce58b81e41 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/9] 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 99ad5a4a2ace4f76e4011ae44d890a958899b8b8 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/9] 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 ddc42ed3cc1307e286eebae54aaab4fc7df39823 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/9] 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 da03864cd6f15..e58cf251c396f 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9948,8 +9948,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 1aa9a46f60fbf370d3ed5e948c1b6cc1a68d5f1e 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/9] 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 9d874d9421a97cf2b54c65891476cffbd072af34 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/9] 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");
>From 0583aac5aa19b4aed535848d3d1f47ac63a3ef38 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 17 Feb 2026 01:18:11 -0500
Subject: [PATCH 7/9] check-refined
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 6 ------
1 file changed, 6 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 0098b9c64ae3e..9e4fe5de5b87d 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8060,12 +8060,6 @@ class MappableExprsHandler {
assert(StrideExpr->getType()->isIntegerType() &&
"Stride expression must be of integer type");
- // If the stride involves member access or array subscript, it's
- // non-contiguous.
- const Expr *S = StrideExpr->IgnoreParenImpCasts();
- if (isa<MemberExpr>(S) || isa<ArraySubscriptExpr>(S))
- return true;
-
// If stride is not evaluatable as a constant, treat as
// non-contiguous.
const auto Constant =
>From 60133b6f385e4278abfde750d4472b80e13ef08d Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Wed, 18 Feb 2026 05:58:25 -0500
Subject: [PATCH 8/9] simplifying_PR
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 29 +-
.../strided_update_count_expression.c | 132 --------
.../strided_update_count_expression_complex.c | 287 ------------------
.../strided_update_count_expression_misc.c | 93 ------
..._update_multiple_arrays_count_expression.c | 162 ----------
.../target_update_ptr_count_expression.c | 99 ------
...t_update_strided_struct_count_expression.c | 97 ------
7 files changed, 8 insertions(+), 891 deletions(-)
delete mode 100644 offload/test/offloading/strided_update_count_expression.c
delete mode 100644 offload/test/offloading/strided_update_count_expression_complex.c
delete mode 100644 offload/test/offloading/strided_update_count_expression_misc.c
delete mode 100644 offload/test/offloading/strided_update_multiple_arrays_count_expression.c
delete mode 100644 offload/test/offloading/target_update_ptr_count_expression.c
delete mode 100644 offload/test/offloading/target_update_strided_struct_count_expression.c
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index e58cf251c396f..25f4da7c90d90 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9945,29 +9945,16 @@ 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)) {
- ConstSizes[I] = 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;
continue;
}
}
diff --git a/offload/test/offloading/strided_update_count_expression.c b/offload/test/offloading/strided_update_count_expression.c
deleted file mode 100644
index e18fc881a8021..0000000000000
--- a/offload/test/offloading/strided_update_count_expression.c
+++ /dev/null
@@ -1,132 +0,0 @@
-// 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>
-
-void test_1_update_from() {
- int len = 10;
- double data[len];
-
- // Initialize data on host
- 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]);
-}
-
-void test_2_update_to() {
- int len = 10;
- double data[len];
-
- 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]);
-}
-
-int main() {
- test_1_update_from();
- test_2_update_to();
- 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
deleted file mode 100644
index 560c4e767f882..0000000000000
--- a/offload/test/offloading/strided_update_count_expression_complex.c
+++ /dev/null
@@ -1,287 +0,0 @@
-// 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];
-};
-
-void test_1_complex_count_expressions() {
- struct Data s1, s2;
- 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 - 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
- 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]);
-}
-
-void test_2_complex_count_with_offset() {
- struct Data s1, s2;
- 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 - 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
- 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]);
-}
-
-// 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
-
-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
deleted file mode 100644
index 6c13eae4d37e2..0000000000000
--- a/offload/test/offloading/strided_update_count_expression_misc.c
+++ /dev/null
@@ -1,93 +0,0 @@
-// 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>
-
-void test_1_modulo_count() {
- int len1 = 10;
- int divisor = 5;
- double data1[len1];
-
- // 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])
- {
-#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]);
-}
-
-void test_2_large_stride_count() {
- int len2 = 10;
- int stride = 5;
- double data2[len2];
-
- // 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])
- {
-#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]);
-}
-
-int main() {
- test_1_modulo_count();
- test_2_large_stride_count();
- 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
deleted file mode 100644
index a1472cacc4a39..0000000000000
--- a/offload/test/offloading/strided_update_multiple_arrays_count_expression.c
+++ /dev/null
@@ -1,162 +0,0 @@
-// 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>
-
-void test_1_update_from_multiple() {
- int n1 = 10, n2 = 10;
- double arr1[n1], arr2[n2];
-
-#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]);
-}
-
-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;
- }
- 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]);
-}
-
-int main() {
- test_1_update_from_multiple();
- test_2_update_to_multiple();
- 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/target_update_ptr_count_expression.c b/offload/test/offloading/target_update_ptr_count_expression.c
deleted file mode 100644
index c4b9fd566d401..0000000000000
--- a/offload/test/offloading/target_update_ptr_count_expression.c
+++ /dev/null
@@ -1,99 +0,0 @@
-// 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_strided_struct_count_expression.c b/offload/test/offloading/target_update_strided_struct_count_expression.c
deleted file mode 100644
index 1c1fd005c405f..0000000000000
--- a/offload/test/offloading/target_update_strided_struct_count_expression.c
+++ /dev/null
@@ -1,97 +0,0 @@
-// 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
>From 3be98c72ce49494d774978e566e15ff9ff4c88d8 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 2 Mar 2026 03:18:51 -0500
Subject: [PATCH 9/9] codegen test added
---
.../target_update_variable_stride_codegen.c | 49 +++++++++++++++++++
1 file changed, 49 insertions(+)
create mode 100644 clang/test/OpenMP/target_update_variable_stride_codegen.c
diff --git a/clang/test/OpenMP/target_update_variable_stride_codegen.c b/clang/test/OpenMP/target_update_variable_stride_codegen.c
new file mode 100644
index 0000000000000..6a4b646e8350b
--- /dev/null
+++ b/clang/test/OpenMP/target_update_variable_stride_codegen.c
@@ -0,0 +1,49 @@
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK --check-prefix CK-64
+// RUN: %clang_cc1 -DCK -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK --check-prefix CK-32
+
+// RUN: %clang_cc1 -DCK -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK
+
+// Test that variable stride expressions in target update directives correctly
+// set the OMP_MAP_NON_CONTIG flag (0x100000000000) in the generated IR.
+// OMP_MAP_TO = 0x1, OMP_MAP_FROM = 0x2
+// NON_CONTIG | TO = 0x100000000001, NON_CONTIG | FROM = 0x100000000002
+
+extern int get_stride();
+
+// CK-64-DAG: [[MTYPE_VAR_STRIDE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x100000000001]]]
+// CK-32-DAG: [[MTYPE_VAR_STRIDE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x100000000001]]]
+// CK-64-DAG: [[MTYPE_VAR_STRIDE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x100000000002]]]
+// CK-32-DAG: [[MTYPE_VAR_STRIDE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x100000000002]]]
+// CK-64-DAG: [[MTYPE_CONST_STRIDE_ONE:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1]]]
+// CK-32-DAG: [[MTYPE_CONST_STRIDE_ONE:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1]]]
+
+void test_variable_stride_to() {
+ int stride = get_stride();
+ int data[10];
+ // CK-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr {{%.+}}, ptr {{%.+}}, ptr @.offload_sizes, ptr [[MTYPE_VAR_STRIDE_TO]], ptr null, ptr null)
+ #pragma omp target update to(data[0:5:stride])
+}
+
+void test_variable_stride_from() {
+ int stride = get_stride();
+ int data[10];
+ // CK-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr {{%.+}}, ptr {{%.+}}, ptr @.offload_sizes{{.*}}, ptr [[MTYPE_VAR_STRIDE_FROM]], ptr null, ptr null)
+ #pragma omp target update from(data[0:5:stride])
+}
+
+void test_constant_stride_one() {
+ int data[10];
+ // CK-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr {{%.+}}, ptr {{%.+}}, ptr @.offload_sizes{{.*}}, ptr [[MTYPE_CONST_STRIDE_ONE]], ptr null, ptr null)
+ #pragma omp target update to(data[0:5:1])
+}
+
+#endif // CK
+#endif // HEADER
More information about the cfe-commits
mailing list