[clang] [llvm] [OpenMP][Offload] Fix base pointer handling for strided target update in structs with pointer members (PR #176914)
Amit Tiwari via cfe-commits
cfe-commits at lists.llvm.org
Tue Feb 17 01:27:39 PST 2026
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/176914
>From 6f610bd8d4f6b149b07a1a4e7bfa66100cfcbbbd 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/8] expression_semantics_patch
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 29 ++++++++++++++++-------
1 file changed, 21 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 8d7a207a91f5a..578e63351ca5f 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9773,16 +9773,29 @@ Error OpenMPIRBuilder::emitOffloadingArrays(
ConstantInt::get(Int64Ty, 0));
SmallBitVector RuntimeSizes(CombinedInfo.Sizes.size());
for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) {
+ bool IsNonContigEntry =
+ IsNonContiguous &&
+ (static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ CombinedInfo.Types[I] &
+ OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0);
+ // For NON_CONTIG entries ArgSizes must carry the dimension count
+ // (number of descriptor_dim records) – NOT the byte size expression.
+ // Variable subsection forms (e.g. 0:s.len/2:2) previously produced a
+ // non-constant size so we marked them runtime and stored the byte size,
+ // leading the runtime to treat it as DimSize and overrun descriptors.
+ if (IsNonContigEntry) {
+ // Dims must be long enough and positive.
+ assert(I < CombinedInfo.NonContigInfo.Dims.size() &&
+ "Induction variable is in-bounds with the NON_CONTIG Dims array");
+ const uint64_t DimCount = CombinedInfo.NonContigInfo.Dims[I];
+ assert(DimCount > 0 && "NON_CONTIG DimCount must be > 0");
+ ConstSizes[I] =
+ ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]);
+ continue;
+ }
if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) {
if (!isa<ConstantExpr>(CI) && !isa<GlobalValue>(CI)) {
- if (IsNonContiguous &&
- static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- CombinedInfo.Types[I] &
- OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG))
- ConstSizes[I] =
- ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]);
- else
- ConstSizes[I] = CI;
+ ConstSizes[I] = CI;
continue;
}
}
>From 6f16202c8a14ac3125a38a47be83f41ea54bb2f3 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/8] variable_stride_fix
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 16 ++++++++++++++++
1 file changed, 16 insertions(+)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 8981a0de6d0e4..027d1fb26bc97 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8033,12 +8033,28 @@ class MappableExprsHandler {
const Expr *StrideExpr = OASE->getStride();
if (!StrideExpr)
return false;
+
+ assert(StrideExpr->getType()->isIntegerType() &&
+ "Stride expression must be of integer type");
+
+ // If the stride is a variable (not a constant), it's non-contiguous.
+ const Expr *S = StrideExpr->IgnoreParenImpCasts();
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(S)) {
+ if (isa<VarDecl>(DRE->getDecl()) ||
+ isa<ParmVarDecl>(DRE->getDecl()))
+ return true;
+ }
+ if (isa<MemberExpr>(S) || isa<ArraySubscriptExpr>(S))
+ return true;
+ // If stride is not evaluatable as a constant, treat as
+ // non-contiguous.
const auto Constant =
StrideExpr->getIntegerConstantExpr(CGF.getContext());
if (!Constant)
return false;
+ // Treat non-unitary strides as non-contiguous.
return !Constant->isOne();
});
>From daebaaead7f5b432d63e677d49f0c80ba8abe491 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 20 Jan 2026 06:15:41 -0500
Subject: [PATCH 3/8] fix base pointer alignment
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 6 ++++--
1 file changed, 4 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 027d1fb26bc97..f4ae2d7d261bc 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8299,7 +8299,9 @@ class MappableExprsHandler {
(Next == CE && MapType != OMPC_MAP_unknown)) {
if (!IsMappingWholeStruct) {
CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
- CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF));
+ CombinedInfo.BasePointers.push_back(IsNonContiguous
+ ? LB.emitRawPointer(CGF)
+ : BP.emitRawPointer(CGF));
CombinedInfo.DevicePtrDecls.push_back(nullptr);
CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
CombinedInfo.Pointers.push_back(LB.emitRawPointer(CGF));
@@ -8407,7 +8409,7 @@ class MappableExprsHandler {
break;
// The pointer becomes the base for the next element.
- if (Next != CE)
+ if (Next != CE && !IsNonContiguous)
BP = IsMemberReference ? LowestElem : LB;
if (!IsPartialMapped)
IsExpressionFirstInfo = false;
>From 761cbba6df6575dc4f53e264f40ffb2299de5ad0 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 20 Jan 2026 06:16:50 -0500
Subject: [PATCH 4/8] tests added
---
..._update_strided_struct_ptr_messages_from.c | 40 +++++++++
...trided_struct_ptr_multiple_messages_from.c | 47 ++++++++++
...strided_struct_ptr_partial_messages_from.c | 32 +++++++
.../target_update_strided_struct_ptr_from.c | 87 +++++++++++++++++++
..._update_strided_struct_ptr_multiple_from.c | 81 +++++++++++++++++
...t_update_strided_struct_ptr_partial_from.c | 67 ++++++++++++++
6 files changed, 354 insertions(+)
create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c
create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_from.c
create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c
create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_partial_from.c
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c
new file mode 100644
index 0000000000000..d86ce9e89766b
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+#define N 16
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = N;
+ s.data = (double *)__builtin_alloca(N * sizeof(double));
+
+ // Valid strided array sections with pointer member
+ #pragma omp target update from(s.data[0:4:2]) // OK
+ {}
+
+ #pragma omp target update from(s.data[1:3:2]) // OK
+ {}
+
+ // Missing stride (default = 1)
+ #pragma omp target update from(s.data[0:4]) // OK
+ {}
+
+ // Invalid stride expressions
+ #pragma omp target update from(s.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ #pragma omp target update from(s.data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Missing colon
+ #pragma omp target update from(s.data[0:4 2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+
+ // Too many colons
+ #pragma omp target update from(s.data[0:4:2:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c
new file mode 100644
index 0000000000000..7020ccb77d231
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c
@@ -0,0 +1,47 @@
+// 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;
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s1, s2;
+ s1.len = N;
+ s1.data = (double *)__builtin_alloca(N * sizeof(double));
+ s2.len = N;
+ s2.data = (double *)__builtin_alloca(N * sizeof(double));
+
+ // Multiple valid strided updates
+ #pragma omp target update from(s1.data[0:10:2], s2.data[0:7:3]) // OK
+ {}
+
+ // Mixed: one with stride, one without
+ #pragma omp target update from(s1.data[0:N], s2.data[0:5:2]) // OK
+ {}
+
+ int stride1 = 2;
+ int stride2 = 3;
+
+ // Multiple with expression strides
+ #pragma omp target update from(s1.data[1:5:stride1], s2.data[0:4:stride2]) // OK
+ {}
+
+ // One valid, one invalid
+ #pragma omp target update from(s1.data[0:5:2], s2.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+
+ #pragma omp target update from(s1.data[0:5:-1], s2.data[0:4:2]) // expected-error {{section stride is evaluated to a non-positive value -1}}
+
+ #pragma omp target update from(s1.data[0:5:0], s2.data[0:4:1]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+
+ // Syntax errors
+ #pragma omp target update from(s1.data[0:5:2], s2.data[0:4 3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+ {}
+
+ #pragma omp target update from(s1.data[0:5:2:3], s2.data[0:4:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+ {}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c
new file mode 100644
index 0000000000000..4c835d3bef6f0
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+// expected-no-diagnostics
+
+#define N 24
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = N;
+ s.data = (double *)__builtin_alloca(N * sizeof(double));
+
+ // Valid partial strided updates with pointer member
+ #pragma omp target update from(s.data[0:2:10]) // OK - partial coverage
+ {}
+
+ // Stride larger than length
+ #pragma omp target update from(s.data[0:2:20]) // OK
+ {}
+
+ // Valid: complex expressions
+ int offset = 1;
+
+ // Runtime-dependent stride expressions
+ #pragma omp target update from(s.data[0:4:offset+1]) // OK
+ {}
+
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_strided_struct_ptr_from.c b/offload/test/offloading/target_update_strided_struct_ptr_from.c
new file mode 100644
index 0000000000000..9785e494ed2ec
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_ptr_from.c
@@ -0,0 +1,87 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that #pragma omp target update from(s.data[0:s.len/2:2])
+// correctly updates every second element (stride 2) from the device to the host
+// using a struct with pointer-to-array member.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16
+
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len])
+
+int main() {
+ T s;
+ s.len = N;
+ s.data = (double *)calloc(N, sizeof(double));
+
+ printf("original host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f\n", s.data[i]);
+ printf("\n");
+
+#pragma omp target data map(mapper(custom), to : s)
+ {
+// Execute on device - modify even-indexed elements
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i += 2) {
+ s.data[i] = 10.0;
+ }
+ }
+
+// Update only even indices (0,2,4,6,8,10,12,14) - s.len/2 elements with stride
+// 2
+#pragma omp target update from(s.data[0 : s.len / 2 : 2])
+ }
+
+ printf("device array values after update from:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f\n", s.data[i]);
+ printf("\n");
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 0.0
+
+ // CHECK: device array values after update from:
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 0.0
+
+ free(s.data);
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c b/offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c
new file mode 100644
index 0000000000000..6ad84c39b717b
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c
@@ -0,0 +1,81 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that multiple strided target updates work correctly
+// with struct containing pointer-to-array member.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 12
+
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len])
+
+int main() {
+ T s1, s2;
+ s1.len = N;
+ s1.data = (double *)calloc(N, sizeof(double));
+ s2.len = N;
+ s2.data = (double *)calloc(N, sizeof(double));
+
+ printf("original s1 values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s1.data[i]);
+ printf("\n");
+
+ printf("original s2 values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s2.data[i]);
+ printf("\n\n");
+
+#pragma omp target data map(mapper(custom), to : s1, s2)
+ {
+// Modify on device
+#pragma omp target
+ {
+ // s1: set even indices to 10
+ for (int i = 0; i < s1.len; i += 2) {
+ s1.data[i] = 10.0;
+ }
+ // s2: set multiples of 3 to 10
+ for (int i = 0; i < s2.len; i += 3) {
+ s2.data[i] = 10.0;
+ }
+ }
+
+// Multiple strided updates: s1 even (s1.len/2 elements, stride 2), s2 multiples
+// of 3 (s2.len/3 elements, stride 3)
+#pragma omp target update from(s1.data[0 : s1.len / 2 : 2], \
+ s2.data[0 : s2.len / 3 : 3])
+ }
+
+ printf("s1 after update (even indices):\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s1.data[i]);
+ printf("\n");
+
+ printf("s2 after update (multiples of 3):\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s2.data[i]);
+ printf("\n");
+
+ // CHECK: original s1 values:
+ // CHECK-NEXT: 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
+
+ // CHECK: original s2 values:
+ // CHECK-NEXT: 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
+
+ // CHECK: s1 after update (even indices):
+ // CHECK-NEXT: 10.0 0.0 10.0 0.0 10.0 0.0 10.0 0.0 10.0 0.0 10.0 0.0
+
+ // CHECK: s2 after update (multiples of 3):
+ // CHECK-NEXT: 10.0 0.0 0.0 10.0 0.0 0.0 10.0 0.0 0.0 10.0 0.0 0.0
+
+ free(s1.data);
+ free(s2.data);
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_strided_struct_ptr_partial_from.c b/offload/test/offloading/target_update_strided_struct_ptr_partial_from.c
new file mode 100644
index 0000000000000..192a8caf7c125
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_ptr_partial_from.c
@@ -0,0 +1,67 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that #pragma omp target update from(s.data[0:N/5:5])
+// correctly updates partial strided elements (stride larger than update count)
+// from device to host using a struct with pointer-to-array member.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 20
+
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len])
+
+int main() {
+ T s;
+ s.len = N;
+ s.data = (double *)calloc(N, sizeof(double));
+
+ printf("original host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s.data[i]);
+ printf("\n\n");
+
+#pragma omp target data map(mapper(custom), tofrom : s)
+ {
+// Set all elements to 20 on device
+#pragma omp target map(mapper(custom), tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = 20.0; // Set all to 20 on device
+ }
+ }
+
+// Modify specific elements on device (only first 4 stride positions)
+#pragma omp target map(mapper(custom), tofrom : s)
+ {
+ s.data[0] = 10.0;
+ s.data[5] = 10.0;
+ s.data[10] = 10.0;
+ s.data[15] = 10.0;
+ }
+
+// Update indices 0, 5, 10, 15 only (N/5 = 4 elements with stride 5)
+#pragma omp target update from(s.data[0 : N / 5 : 5])
+ }
+
+ printf("device array values after partial stride update:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s.data[i]);
+ printf("\n");
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0
+ // 0.0 0.0 0.0 0.0
+
+ // CHECK: device array values after partial stride update:
+ // CHECK-NEXT: 10.0 0.0 0.0 0.0 0.0 10.0 0.0 0.0 0.0 0.0 10.0 0.0 0.0 0.0
+ // 0.0 10.0 0.0 0.0 0.0 0.0
+
+ free(s.data);
+ return 0;
+}
>From 487cdedc995d72029954fe34407d06b769bd49e0 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 20 Jan 2026 06:18:41 -0500
Subject: [PATCH 5/8] refined pointer fix
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 4 +---
1 file changed, 1 insertion(+), 3 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index f4ae2d7d261bc..acb429c3ce7b3 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8299,9 +8299,7 @@ class MappableExprsHandler {
(Next == CE && MapType != OMPC_MAP_unknown)) {
if (!IsMappingWholeStruct) {
CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
- CombinedInfo.BasePointers.push_back(IsNonContiguous
- ? LB.emitRawPointer(CGF)
- : BP.emitRawPointer(CGF));
+ CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF));
CombinedInfo.DevicePtrDecls.push_back(nullptr);
CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
CombinedInfo.Pointers.push_back(LB.emitRawPointer(CGF));
>From 75f49b84d69f8be7c3fa4ea803782f853d667cbc Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 20 Jan 2026 06:26:13 -0500
Subject: [PATCH 6/8] all tests covered
---
...rget_update_strided_struct_ptr_codegen.cpp | 73 ++++++++++++
...et_update_strided_struct_ptr_messages_to.c | 40 +++++++
..._strided_struct_ptr_multiple_messages_to.c | 47 ++++++++
...e_strided_struct_ptr_partial_messages_to.c | 32 ++++++
...et_update_strided_struct_ptr_multiple_to.c | 101 +++++++++++++++++
...get_update_strided_struct_ptr_partial_to.c | 76 +++++++++++++
.../target_update_strided_struct_ptr_to.c | 106 ++++++++++++++++++
7 files changed, 475 insertions(+)
create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp
create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c
create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c
create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_partial_to.c
create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_to.c
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp b/clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp
new file mode 100644
index 0000000000000..a2e5978a81eb7
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp
@@ -0,0 +1,73 @@
+// Test codegen for strided target update with struct containing pointer-to-array member
+// RUN: %clang_cc1 -DCK27 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK27
+// RUN: %clang_cc1 -DCK27 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK27
+
+// RUN: %clang_cc1 -DCK27 -verify -Wno-vla -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 -DCK27 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+// expected-no-diagnostics
+
+#ifdef CK27
+#ifndef CK27_INCLUDED
+#define CK27_INCLUDED
+
+// Verify that non-contiguous map type flag is set (bit 48)
+// 17592186044418 = 0x1000000000002 (OMP_MAP_NON_CONTIG | OMP_MAP_FROM)
+// 17592186044417 = 0x1000000000001 (OMP_MAP_NON_CONTIG | OMP_MAP_TO)
+// CK27-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 17592186044418]
+// CK27-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 17592186044417]
+
+struct T {
+ double *data;
+ int len;
+};
+
+// CK27-LABEL: define {{.*}}void @{{.*}}test_strided_update_from{{.*}}(
+void test_strided_update_from(int arg) {
+ T s;
+ s.len = 16;
+ s.data = new double[16];
+
+ for (int i = 0; i < 16; i++) {
+ s.data[i] = i;
+ }
+
+ // Verify the stride descriptor is created with correct values:
+ // - offset = 0
+ // - count = 4 (number of elements to update)
+ // - stride = 16 (2 * sizeof(double) = 2 * 8 = 16 bytes)
+ // CK27-DAG: store i64 0, ptr %{{.+}}, align 8
+ // CK27-DAG: store i64 4, ptr %{{.+}}, align 8
+ // CK27-DAG: store i64 16, ptr %{{.+}}, align 8
+
+ // Verify __tgt_target_data_update_mapper is called
+ // CK27: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 {{1|signext 1}}, ptr %{{.+}}, ptr %{{.+}}, ptr @{{.+}}, ptr @.offload_maptypes{{.*}}, ptr null, ptr null)
+
+ #pragma omp target update from(s.data[0:4:2])
+
+ delete[] s.data;
+}
+
+// CK27-LABEL: define {{.*}}void @{{.*}}test_strided_update_to{{.*}}(
+void test_strided_update_to(int arg) {
+ T s;
+ s.len = 16;
+ s.data = new double[16];
+
+ for (int i = 0; i < 16; i++) {
+ s.data[i] = i;
+ }
+
+ // Verify __tgt_target_data_update_mapper is called with TO map type
+ // CK27: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 {{1|signext 1}}, ptr %{{.+}}, ptr %{{.+}}, ptr @{{.+}}, ptr @.offload_maptypes{{.*}}, ptr null, ptr null)
+
+ #pragma omp target update to(s.data[0:4:2])
+
+ delete[] s.data;
+}
+
+#endif // CK27_INCLUDED
+#endif // CK27
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c
new file mode 100644
index 0000000000000..012f484a6ae36
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+#define N 16
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = N;
+ s.data = (double *)__builtin_alloca(N * sizeof(double));
+
+ // Valid strided array sections with pointer member
+ #pragma omp target update to(s.data[0:4:2]) // OK
+ {}
+
+ #pragma omp target update to(s.data[1:3:2]) // OK
+ {}
+
+ // Missing stride (default = 1)
+ #pragma omp target update to(s.data[0:4]) // OK
+ {}
+
+ // Invalid stride expressions
+ #pragma omp target update to(s.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ #pragma omp target update to(s.data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Missing colon
+ #pragma omp target update to(s.data[0:4 2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+
+ // Too many colons
+ #pragma omp target update to(s.data[0:4:2:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c
new file mode 100644
index 0000000000000..05278308ad173
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c
@@ -0,0 +1,47 @@
+// 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;
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s1, s2;
+ s1.len = N;
+ s1.data = (double *)__builtin_alloca(N * sizeof(double));
+ s2.len = N;
+ s2.data = (double *)__builtin_alloca(N * sizeof(double));
+
+ // Multiple valid strided updates (to clause)
+ #pragma omp target update to(s1.data[0:10:2], s2.data[0:7:3]) // OK
+ {}
+
+ // Mixed: one with stride, one without
+ #pragma omp target update to(s1.data[0:N], s2.data[0:5:2]) // OK
+ {}
+
+ int stride1 = 2;
+ int stride2 = 3;
+
+ // Multiple with expression strides
+ #pragma omp target update to(s1.data[1:5:stride1], s2.data[0:4:stride2]) // OK
+ {}
+
+ // One valid, one invalid
+ #pragma omp target update to(s1.data[0:5:2], s2.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+
+ #pragma omp target update to(s1.data[0:5:-1], s2.data[0:4:2]) // expected-error {{section stride is evaluated to a non-positive value -1}}
+
+ #pragma omp target update to(s1.data[0:5:0], s2.data[0:4:1]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+
+ // Syntax errors
+ #pragma omp target update to(s1.data[0:5:2], s2.data[0:4 3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+ {}
+
+ #pragma omp target update to(s1.data[0:5:2:3], s2.data[0:4:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+ {}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c
new file mode 100644
index 0000000000000..d62a6c640d0b3
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+// expected-no-diagnostics
+
+#define N 24
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = N;
+ s.data = (double *)__builtin_alloca(N * sizeof(double));
+
+ // Valid partial strided updates with pointer member (to clause)
+ #pragma omp target update to(s.data[0:2:10]) // OK - partial coverage
+ {}
+
+ // Stride larger than length
+ #pragma omp target update to(s.data[0:2:20]) // OK
+ {}
+
+ // Valid: complex expressions
+ int offset = 1;
+
+ // Runtime-dependent stride expressions
+ #pragma omp target update to(s.data[0:4:offset+1]) // OK
+ {}
+
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c b/offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c
new file mode 100644
index 0000000000000..267c9f1db1b9f
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c
@@ -0,0 +1,101 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that multiple strided target updates to device work
+// correctly with struct containing pointer-to-array member.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 12
+
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+#pragma omp declare mapper(custom : T v) \
+ map(tofrom : v, v.len, v.data[0 : v.len])
+
+int main() {
+ T s1, s2;
+ s1.len = N;
+ s1.data = (double *)calloc(N, sizeof(double));
+ s2.len = N;
+ s2.data = (double *)calloc(N, sizeof(double));
+
+ // Initialize structs on host
+ for (int i = 0; i < N; i++) {
+ s1.data[i] = i;
+ s2.data[i] = i;
+ }
+
+ printf("original s1 values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s1.data[i]);
+ printf("\n");
+
+ printf("original s2 values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s2.data[i]);
+ printf("\n\n");
+
+#pragma omp target data map(tofrom : s1, s2)
+ {
+// Initialize device struct arrays to 20
+#pragma omp target
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.data[i] = 20.0;
+ s2.data[i] = 20.0;
+ }
+ }
+
+ // Modify host: s1 even indices, s2 multiples of 3
+ for (int i = 0; i < s1.len; i += 2) {
+ s1.data[i] = 10.0;
+ }
+ for (int i = 0; i < s2.len; i += 3) {
+ s2.data[i] = 10.0;
+ }
+
+// Multiple strided updates to device: s1 even (s1.len/2 elements, stride 2), s2
+// multiples of 3 (s2.len/3 elements, stride 3)
+#pragma omp target update to(s1.data[0 : s1.len / 2 : 2], \
+ s2.data[0 : s2.len / 3 : 3])
+
+// Verify update on device by adding 5
+#pragma omp target
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.data[i] += 5.0;
+ s2.data[i] += 5.0;
+ }
+ }
+ } // Exit target data - tofrom mapper copies data back
+
+ printf("s1 after update to device (even indices):\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s1.data[i]);
+ printf("\n");
+
+ printf("s2 after update to device (multiples of 3):\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s2.data[i]);
+ printf("\n");
+
+ // CHECK: original s1 values:
+ // CHECK-NEXT: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0
+
+ // CHECK: original s2 values:
+ // CHECK-NEXT: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0
+
+ // CHECK: s1 after update to device (even indices):
+ // CHECK-NEXT: 15.0 25.0 15.0 25.0 15.0 25.0 15.0 25.0 15.0 25.0 15.0 25.0
+
+ // CHECK: s2 after update to device (multiples of 3):
+ // CHECK-NEXT: 15.0 25.0 25.0 15.0 25.0 25.0 15.0 25.0 25.0 15.0 25.0 25.0
+
+ free(s1.data);
+ free(s2.data);
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_strided_struct_ptr_partial_to.c b/offload/test/offloading/target_update_strided_struct_ptr_partial_to.c
new file mode 100644
index 0000000000000..35772f7c642e4
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_ptr_partial_to.c
@@ -0,0 +1,76 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that #pragma omp target update to(s.data[0:N/5:5]) correctly
+// updates partial strided elements (stride larger than update count) from host
+// to device using a struct with pointer-to-array member.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 20
+
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+#pragma omp declare mapper(custom : T v) \
+ map(tofrom : v, v.len, v.data[0 : v.len])
+
+int main() {
+ T s;
+ s.len = N;
+ s.data = (double *)calloc(N, sizeof(double));
+
+ // Initialize struct data on host
+ for (int i = 0; i < N; i++) {
+ s.data[i] = i;
+ }
+
+ printf("original host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s.data[i]);
+ printf("\n\n");
+
+#pragma omp target data map(tofrom : s)
+ {
+// Initialize device struct arrays to 20
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = 20.0;
+ }
+ }
+
+ // Modify host elements: indices 0, 5, 10, 15 only
+ s.data[0] = 10.0;
+ s.data[5] = 10.0;
+ s.data[10] = 10.0;
+ s.data[15] = 10.0;
+
+// Update indices 0, 5, 10, 15 only (N/5 = 4 elements with stride 5) to device
+#pragma omp target update to(s.data[0 : N / 5 : 5])
+
+// Execute on device - add 5 to verify update worked
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += 5.0;
+ }
+ }
+ } // Exit target data - tofrom mapper copies data back
+
+ printf("device array values after partial stride update to:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f ", s.data[i]);
+ printf("\n");
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0 12.0 13.0 14.0 15.0 16.0 17.0 18.0 19.0
+
+ // CHECK: device array values after partial stride update to:
+ // CHECK-NEXT: 15.0 25.0 25.0 25.0 25.0 15.0 25.0 25.0 25.0 25.0 15.0 25.0 25.0 25.0 25.0 15.0 25.0 25.0 25.0 25.0
+
+ free(s.data);
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_strided_struct_ptr_to.c b/offload/test/offloading/target_update_strided_struct_ptr_to.c
new file mode 100644
index 0000000000000..690802e202065
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_ptr_to.c
@@ -0,0 +1,106 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that #pragma omp target update to(s.data[0:s.len/2:2])
+// correctly updates every second element (stride 2) from the host to the device
+// using a struct with pointer-to-array member.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16
+
+typedef struct {
+ double *data;
+ int len;
+} T;
+
+#pragma omp declare mapper(custom : T v) \
+ map(tofrom : v, v.len, v.data[0 : v.len])
+
+int main() {
+ T s;
+ s.len = N;
+ s.data = (double *)calloc(N, sizeof(double));
+
+ // Initialize struct data on host
+ for (int i = 0; i < N; i++) {
+ s.data[i] = i;
+ }
+
+ printf("original host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f\n", s.data[i]);
+ printf("\n");
+
+#pragma omp target data map(tofrom : s)
+ {
+// Set device data to 20
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = 20.0;
+ }
+ }
+
+ // Modify host even-indexed elements
+ for (int i = 0; i < N; i += 2) {
+ s.data[i] = 10.0;
+ }
+
+// Update only even indices (0,2,4,6,8,10,12,14) to device - s.len/2 elements
+// with stride 2
+#pragma omp target update to(s.data[0 : s.len / 2 : 2])
+
+// Execute on device - add 5 to verify update worked
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += 5.0;
+ }
+ }
+ } // Exit target data - tofrom mapper copies data back
+
+ printf("device array values after update to:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f\n", s.data[i]);
+ printf("\n");
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 11.0
+ // CHECK-NEXT: 12.0
+ // CHECK-NEXT: 13.0
+ // CHECK-NEXT: 14.0
+ // CHECK-NEXT: 15.0
+
+ // CHECK: device array values after update to:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+
+ free(s.data);
+ return 0;
+}
>From 66ad400a86a79c1a7edf7d78d0e28490a42809b6 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 20 Jan 2026 07:44:30 -0500
Subject: [PATCH 7/8] formatted
---
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 acb429c3ce7b3..1e006fb670021 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8033,9 +8033,9 @@ class MappableExprsHandler {
const Expr *StrideExpr = OASE->getStride();
if (!StrideExpr)
return false;
-
+
assert(StrideExpr->getType()->isIntegerType() &&
- "Stride expression must be of integer type");
+ "Stride expression must be of integer type");
// If the stride is a variable (not a constant), it's non-contiguous.
const Expr *S = StrideExpr->IgnoreParenImpCasts();
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 578e63351ca5f..6f44a9649abeb 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9776,8 +9776,8 @@ Error OpenMPIRBuilder::emitOffloadingArrays(
bool IsNonContigEntry =
IsNonContiguous &&
(static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- CombinedInfo.Types[I] &
- OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0);
+ CombinedInfo.Types[I] &
+ OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0);
// For NON_CONTIG entries ArgSizes must carry the dimension count
// (number of descriptor_dim records) – NOT the byte size expression.
// Variable subsection forms (e.g. 0:s.len/2:2) previously produced a
>From 26f0615b2827f53fb26f9a3b874985159b42cfd8 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 17 Feb 2026 04:27:14 -0500
Subject: [PATCH 8/8] removed_parsing_checks
---
..._update_strided_struct_ptr_messages_from.c | 40 ----------------
...et_update_strided_struct_ptr_messages_to.c | 40 ----------------
...trided_struct_ptr_multiple_messages_from.c | 47 -------------------
..._strided_struct_ptr_multiple_messages_to.c | 47 -------------------
...strided_struct_ptr_partial_messages_from.c | 32 -------------
...e_strided_struct_ptr_partial_messages_to.c | 32 -------------
6 files changed, 238 deletions(-)
delete mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c
delete mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c
delete mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c
delete mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c
delete mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c
delete mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c
deleted file mode 100644
index d86ce9e89766b..0000000000000
--- a/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c
+++ /dev/null
@@ -1,40 +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 16
-typedef struct {
- double *data;
- int len;
-} T;
-
-int main(int argc, char **argv) {
- T s;
- s.len = N;
- s.data = (double *)__builtin_alloca(N * sizeof(double));
-
- // Valid strided array sections with pointer member
- #pragma omp target update from(s.data[0:4:2]) // OK
- {}
-
- #pragma omp target update from(s.data[1:3:2]) // OK
- {}
-
- // Missing stride (default = 1)
- #pragma omp target update from(s.data[0:4]) // OK
- {}
-
- // Invalid stride expressions
- #pragma omp target update from(s.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
-
- #pragma omp target update from(s.data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
-
- // Missing colon
- #pragma omp target update from(s.data[0:4 2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
- {}
-
- // Too many colons
- #pragma omp target update from(s.data[0:4:2:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
- {}
-
- return 0;
-}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c
deleted file mode 100644
index 012f484a6ae36..0000000000000
--- a/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c
+++ /dev/null
@@ -1,40 +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 16
-typedef struct {
- double *data;
- int len;
-} T;
-
-int main(int argc, char **argv) {
- T s;
- s.len = N;
- s.data = (double *)__builtin_alloca(N * sizeof(double));
-
- // Valid strided array sections with pointer member
- #pragma omp target update to(s.data[0:4:2]) // OK
- {}
-
- #pragma omp target update to(s.data[1:3:2]) // OK
- {}
-
- // Missing stride (default = 1)
- #pragma omp target update to(s.data[0:4]) // OK
- {}
-
- // Invalid stride expressions
- #pragma omp target update to(s.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
-
- #pragma omp target update to(s.data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
-
- // Missing colon
- #pragma omp target update to(s.data[0:4 2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
- {}
-
- // Too many colons
- #pragma omp target update to(s.data[0:4:2:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
- {}
-
- return 0;
-}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c
deleted file mode 100644
index 7020ccb77d231..0000000000000
--- a/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c
+++ /dev/null
@@ -1,47 +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;
- int len;
-} T;
-
-int main(int argc, char **argv) {
- T s1, s2;
- s1.len = N;
- s1.data = (double *)__builtin_alloca(N * sizeof(double));
- s2.len = N;
- s2.data = (double *)__builtin_alloca(N * sizeof(double));
-
- // Multiple valid strided updates
- #pragma omp target update from(s1.data[0:10:2], s2.data[0:7:3]) // OK
- {}
-
- // Mixed: one with stride, one without
- #pragma omp target update from(s1.data[0:N], s2.data[0:5:2]) // OK
- {}
-
- int stride1 = 2;
- int stride2 = 3;
-
- // Multiple with expression strides
- #pragma omp target update from(s1.data[1:5:stride1], s2.data[0:4:stride2]) // OK
- {}
-
- // One valid, one invalid
- #pragma omp target update from(s1.data[0:5:2], s2.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
-
- #pragma omp target update from(s1.data[0:5:-1], s2.data[0:4:2]) // expected-error {{section stride is evaluated to a non-positive value -1}}
-
- #pragma omp target update from(s1.data[0:5:0], s2.data[0:4:1]) // expected-error {{section stride is evaluated to a non-positive value 0}}
-
- // Syntax errors
- #pragma omp target update from(s1.data[0:5:2], s2.data[0:4 3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
- {}
-
- #pragma omp target update from(s1.data[0:5:2:3], s2.data[0:4:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
- {}
-
- return 0;
-}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c
deleted file mode 100644
index 05278308ad173..0000000000000
--- a/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c
+++ /dev/null
@@ -1,47 +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;
- int len;
-} T;
-
-int main(int argc, char **argv) {
- T s1, s2;
- s1.len = N;
- s1.data = (double *)__builtin_alloca(N * sizeof(double));
- s2.len = N;
- s2.data = (double *)__builtin_alloca(N * sizeof(double));
-
- // Multiple valid strided updates (to clause)
- #pragma omp target update to(s1.data[0:10:2], s2.data[0:7:3]) // OK
- {}
-
- // Mixed: one with stride, one without
- #pragma omp target update to(s1.data[0:N], s2.data[0:5:2]) // OK
- {}
-
- int stride1 = 2;
- int stride2 = 3;
-
- // Multiple with expression strides
- #pragma omp target update to(s1.data[1:5:stride1], s2.data[0:4:stride2]) // OK
- {}
-
- // One valid, one invalid
- #pragma omp target update to(s1.data[0:5:2], s2.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
-
- #pragma omp target update to(s1.data[0:5:-1], s2.data[0:4:2]) // expected-error {{section stride is evaluated to a non-positive value -1}}
-
- #pragma omp target update to(s1.data[0:5:0], s2.data[0:4:1]) // expected-error {{section stride is evaluated to a non-positive value 0}}
-
- // Syntax errors
- #pragma omp target update to(s1.data[0:5:2], s2.data[0:4 3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
- {}
-
- #pragma omp target update to(s1.data[0:5:2:3], s2.data[0:4:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
- {}
-
- return 0;
-}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c
deleted file mode 100644
index 4c835d3bef6f0..0000000000000
--- a/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c
+++ /dev/null
@@ -1,32 +0,0 @@
-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-// expected-no-diagnostics
-
-#define N 24
-typedef struct {
- double *data;
- int len;
-} T;
-
-int main(int argc, char **argv) {
- T s;
- s.len = N;
- s.data = (double *)__builtin_alloca(N * sizeof(double));
-
- // Valid partial strided updates with pointer member
- #pragma omp target update from(s.data[0:2:10]) // OK - partial coverage
- {}
-
- // Stride larger than length
- #pragma omp target update from(s.data[0:2:20]) // OK
- {}
-
- // Valid: complex expressions
- int offset = 1;
-
- // Runtime-dependent stride expressions
- #pragma omp target update from(s.data[0:4:offset+1]) // OK
- {}
-
- return 0;
-}
diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c
deleted file mode 100644
index d62a6c640d0b3..0000000000000
--- a/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c
+++ /dev/null
@@ -1,32 +0,0 @@
-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-// expected-no-diagnostics
-
-#define N 24
-typedef struct {
- double *data;
- int len;
-} T;
-
-int main(int argc, char **argv) {
- T s;
- s.len = N;
- s.data = (double *)__builtin_alloca(N * sizeof(double));
-
- // Valid partial strided updates with pointer member (to clause)
- #pragma omp target update to(s.data[0:2:10]) // OK - partial coverage
- {}
-
- // Stride larger than length
- #pragma omp target update to(s.data[0:2:20]) // OK
- {}
-
- // Valid: complex expressions
- int offset = 1;
-
- // Runtime-dependent stride expressions
- #pragma omp target update to(s.data[0:4:offset+1]) // OK
- {}
-
- return 0;
-}
More information about the cfe-commits
mailing list