[clang] [llvm] [Clang][OpenMP] Handled `NonContig` Descriptor `DimCount` (PR #181987)
Amit Tiwari via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 3 01:56:53 PST 2026
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/181987
>From fe3b4e0cbaaebbdc8b5a59590e363d42c4ccf8b0 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/5] 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 2dcfc9a02d196c07901a99599bab26a8a3fa33fb Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Wed, 18 Feb 2026 05:29:00 -0500
Subject: [PATCH 2/5] tests_added
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 4 +-
.../strided_update_count_expression.c | 132 ++++++++
.../strided_update_count_expression_complex.c | 287 ++++++++++++++++++
..._update_multiple_arrays_count_expression.c | 162 ++++++++++
.../target_update_ptr_count_expression.c | 99 ++++++
...t_update_strided_struct_count_expression.c | 97 ++++++
6 files changed, 779 insertions(+), 2 deletions(-)
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_multiple_arrays_count_expression.c
create mode 100644 offload/test/offloading/target_update_ptr_count_expression.c
create 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 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
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..e18fc881a8021
--- /dev/null
+++ b/offload/test/offloading/strided_update_count_expression.c
@@ -0,0 +1,132 @@
+// 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
new file mode 100644
index 0000000000000..560c4e767f882
--- /dev/null
+++ b/offload/test/offloading/strided_update_count_expression_complex.c
@@ -0,0 +1,287 @@
+// 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_multiple_arrays_count_expression.c b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c
new file mode 100644
index 0000000000000..a1472cacc4a39
--- /dev/null
+++ b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c
@@ -0,0 +1,162 @@
+// 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
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_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
>From 2f194562c128f7eb62d2cdf9973519399167a1eb Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Thu, 19 Feb 2026 06:02:56 -0500
Subject: [PATCH 3/5] revised
---
.../ASTMatchers/ASTMatchersNarrowingTest.cpp | 88 +++++++++++++++++++
.../ASTMatchers/ASTMatchersNodeTest.cpp | 46 ++++++++++
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 13 +--
.../Frontend/OpenMPIRBuilderTest.cpp | 48 ++++++++++
...t_update_strided_struct_count_expression.c | 2 +-
5 files changed, 187 insertions(+), 10 deletions(-)
diff --git a/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp b/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp
index b59b93bd9fcd4..d9647c201fc30 100644
--- a/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp
+++ b/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp
@@ -5168,6 +5168,94 @@ TEST_P(ASTMatchersTest, OMPFromClause_DoesNotMatchMapClause) {
EXPECT_TRUE(notMatchesWithOpenMP(Source0, Matcher));
}
+TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_ArraySection_CountExpression) {
+ StringRef Source0 = R"(
+ void foo() {
+ int count = 8;
+ int arr[100];
+ #pragma omp target update from(arr[0:count:2])
+ ;
+ }
+ )";
+
+ auto astUnit =
+ tooling::buildASTFromCodeWithArgs(Source0, {"-fopenmp=libomp"});
+ ASSERT_TRUE(astUnit);
+
+ auto Results = match(ompTargetUpdateDirective().bind("directive"),
+ astUnit->getASTContext());
+ ASSERT_FALSE(Results.empty());
+
+ const auto *Directive =
+ Results[0].getNodeAs<OMPTargetUpdateDirective>("directive");
+ ASSERT_TRUE(Directive);
+
+ OMPFromClause *FromClause = nullptr;
+ for (auto *Clause : Directive->clauses()) {
+ if ((FromClause = dyn_cast<OMPFromClause>(Clause))) {
+ break;
+ }
+ }
+ ASSERT_TRUE(FromClause);
+
+ for (const auto *VarExpr : FromClause->varlist()) {
+ const auto *ArraySection = dyn_cast<ArraySectionExpr>(VarExpr);
+ if (!ArraySection)
+ continue;
+
+ // Verify length expression exists and is not a constant
+ const Expr *Length = ArraySection->getLength();
+ ASSERT_TRUE(Length);
+ EXPECT_FALSE(isa<IntegerLiteral>(Length))
+ << "Expected length to be a variable expression (count)";
+ }
+}
+
+TEST_P(ASTMatchersTest,
+ OMPTargetUpdateDirective_ArraySection_ComplexCountExpression) {
+ StringRef Source0 = R"(
+ void foo() {
+ int len = 16;
+ int count = 8;
+ int arr[100];
+ #pragma omp target update from(arr[0:(len+count)/2:2])
+ ;
+ }
+ )";
+
+ auto astUnit =
+ tooling::buildASTFromCodeWithArgs(Source0, {"-fopenmp=libomp"});
+ ASSERT_TRUE(astUnit);
+
+ auto Results = match(ompTargetUpdateDirective().bind("directive"),
+ astUnit->getASTContext());
+ ASSERT_FALSE(Results.empty());
+
+ const auto *Directive =
+ Results[0].getNodeAs<OMPTargetUpdateDirective>("directive");
+ ASSERT_TRUE(Directive);
+
+ OMPFromClause *FromClause = nullptr;
+ for (auto *Clause : Directive->clauses()) {
+ if ((FromClause = dyn_cast<OMPFromClause>(Clause))) {
+ break;
+ }
+ }
+ ASSERT_TRUE(FromClause);
+
+ for (const auto *VarExpr : FromClause->varlist()) {
+ const auto *ArraySection = dyn_cast<ArraySectionExpr>(VarExpr);
+ if (!ArraySection)
+ continue;
+
+ // Verify length is a complex expression, not a constant
+ const Expr *Length = ArraySection->getLength();
+ ASSERT_TRUE(Length);
+ EXPECT_FALSE(isa<IntegerLiteral>(Length))
+ << "Expected length to be a complex expression ((len+count)/2)";
+ }
+}
+
TEST_P(ASTMatchersTest, HasAnyBase_DirectBase) {
if (!GetParam().isCXX()) {
return;
diff --git a/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp b/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp
index 96f1b5249d859..7338ff5f302f6 100644
--- a/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp
+++ b/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp
@@ -3057,6 +3057,52 @@ void x(int x) {
EXPECT_TRUE(notMatchesWithOpenMP(Source5, Matcher));
}
+TEST(ASTMatchersTestOpenMP, OMPTargetUpdateDirective_CountExpression) {
+ auto Matcher = ompTargetUpdateDirective(hasAnyClause(ompFromClause()));
+
+ StringRef Source0 = R"(
+ void foo() {
+ int len = 16;
+ int arr[100];
+ #pragma omp target update from(arr[0:len/2:2])
+ ;
+ }
+ )";
+ EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
+
+ auto astUnit =
+ tooling::buildASTFromCodeWithArgs(Source0, {"-fopenmp=libomp"});
+ ASSERT_TRUE(astUnit);
+
+ auto Results = match(ompTargetUpdateDirective().bind("directive"),
+ astUnit->getASTContext());
+ ASSERT_FALSE(Results.empty());
+
+ const auto *Directive =
+ Results[0].getNodeAs<OMPTargetUpdateDirective>("directive");
+ ASSERT_TRUE(Directive);
+
+ OMPFromClause *FromClause = nullptr;
+ for (auto *Clause : Directive->clauses()) {
+ if ((FromClause = dyn_cast<OMPFromClause>(Clause))) {
+ break;
+ }
+ }
+ ASSERT_TRUE(FromClause);
+
+ for (const auto *VarExpr : FromClause->varlist()) {
+ const auto *ArraySection = dyn_cast<ArraySectionExpr>(VarExpr);
+ if (!ArraySection)
+ continue;
+
+ const Expr *Length = ArraySection->getLength();
+ ASSERT_TRUE(Length);
+ const auto *LengthLiteral = dyn_cast<IntegerLiteral>(Length);
+ EXPECT_FALSE(LengthLiteral)
+ << "Expected length to be a variable expression";
+ }
+}
+
TEST(ASTMatchersTest, Finder_DynamicOnlyAcceptsSomeMatchers) {
MatchFinder Finder;
EXPECT_TRUE(Finder.addDynamicMatcher(decl(), nullptr));
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index e58cf251c396f..a1d5fa7cff992 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9950,19 +9950,14 @@ Error OpenMPIRBuilder::emitOffloadingArrays(
(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.
+ // For NON_CONTIG entries, ArgSizes stores the dimension count (number of
+ // descriptor_dim records), not the byte size.
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");
+ "Index must be in-bounds for 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]);
+ ConstSizes[I] = ConstantInt::get(Int64Ty, DimCount);
continue;
}
if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) {
diff --git a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
index 003a831833482..c5e9901463a8f 100644
--- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
@@ -7940,4 +7940,52 @@ TEST_F(OpenMPIRBuilderTest, spliceBBWithEmptyBB) {
EXPECT_FALSE(Terminator->getDbgRecordRange().empty());
}
+// Test that NON_CONTIG entries with count expressions store dimension count
+TEST_F(OpenMPIRBuilderTest, EmitOffloadingArraysNonContigCountExpression) {
+ OpenMPIRBuilder OMPBuilder(*M);
+ OMPBuilder.initialize();
+ IRBuilder<> Builder(BB);
+
+ Value *BasePtr = ConstantPointerNull::get(Builder.getPtrTy());
+ Value *CountExpr =
+ Builder.CreateUDiv(Builder.getInt64(16), Builder.getInt64(2));
+ Value *ByteSize =
+ Builder.CreateMul(CountExpr, Builder.getInt64(4)); // Non-constant
+
+ // Set up NON_CONTIG entry with count expression
+ OpenMPIRBuilder::MapInfosTy CombinedInfo;
+ CombinedInfo.BasePointers.push_back(BasePtr);
+ CombinedInfo.Pointers.push_back(BasePtr);
+ CombinedInfo.DevicePointers.push_back(OpenMPIRBuilder::DeviceInfoTy::Pointer);
+ CombinedInfo.Sizes.push_back(ByteSize);
+ CombinedInfo.Types.push_back(static_cast<omp::OpenMPOffloadMappingFlags>(
+ omp::OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG |
+ omp::OpenMPOffloadMappingFlags::OMP_MAP_TO));
+ CombinedInfo.Names.push_back(
+ Builder.CreateGlobalString("data", "data_name", 0, M.get()));
+
+ CombinedInfo.NonContigInfo.IsNonContiguous = true;
+ CombinedInfo.NonContigInfo.Dims.push_back(1);
+ CombinedInfo.NonContigInfo.Offsets.push_back({Builder.getInt64(0)});
+ CombinedInfo.NonContigInfo.Counts.push_back({CountExpr});
+ CombinedInfo.NonContigInfo.Strides.push_back({Builder.getInt64(2)});
+
+ OpenMPIRBuilder::TargetDataInfo Info(true, false);
+ Info.NumberOfPtrs = 1;
+ OpenMPIRBuilder::TargetDataRTArgs RTArgs;
+ using InsertPointTy = OpenMPIRBuilder::InsertPointTy;
+
+ EXPECT_FALSE(OMPBuilder.emitOffloadingArraysAndArgs(
+ InsertPointTy(Builder.saveIP()), InsertPointTy(Builder.saveIP()), Info,
+ RTArgs, CombinedInfo,
+ [](unsigned) -> Expected<Function *> {
+ return static_cast<Function *>(nullptr);
+ },
+ /*IsNonContiguous=*/true));
+
+ // Verify SizesArray is constant, not runtime
+ EXPECT_NE(dyn_cast<GlobalVariable>(RTArgs.SizesArray), nullptr);
+ EXPECT_EQ(dyn_cast<AllocaInst>(RTArgs.SizesArray), nullptr);
+}
+
} // namespace
diff --git a/offload/test/offloading/target_update_strided_struct_count_expression.c b/offload/test/offloading/target_update_strided_struct_count_expression.c
index 1c1fd005c405f..89b0e1b4c7ea7 100644
--- a/offload/test/offloading/target_update_strided_struct_count_expression.c
+++ b/offload/test/offloading/target_update_strided_struct_count_expression.c
@@ -54,7 +54,7 @@ int main() {
}
// Update to device with expression-based count
-#pragma omp target data map(to : s)
+#pragma omp target data map(alloc : s)
{
#pragma omp target update to(s.data[0 : s.len / 2 : 2])
>From 0a5599c753490810094c33c939aa0273c748c223 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Thu, 19 Feb 2026 07:28:52 -0500
Subject: [PATCH 4/5] gpu_set_on
---
llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp | 1 +
1 file changed, 1 insertion(+)
diff --git a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
index c5e9901463a8f..1416f748c3dc8 100644
--- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
@@ -7943,6 +7943,7 @@ TEST_F(OpenMPIRBuilderTest, spliceBBWithEmptyBB) {
// Test that NON_CONTIG entries with count expressions store dimension count
TEST_F(OpenMPIRBuilderTest, EmitOffloadingArraysNonContigCountExpression) {
OpenMPIRBuilder OMPBuilder(*M);
+ OMPBuilder.Config.setIsGPU(true);
OMPBuilder.initialize();
IRBuilder<> Builder(BB);
>From ead2ce4165c01f53e339e19881428c5b67b28aed Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 3 Mar 2026 04:51:24 -0500
Subject: [PATCH 5/5] codegen-test
---
.../target_update_count_expression_codegen.c | 41 +++++++++++++++++++
1 file changed, 41 insertions(+)
create mode 100644 clang/test/OpenMP/target_update_count_expression_codegen.c
diff --git a/clang/test/OpenMP/target_update_count_expression_codegen.c b/clang/test/OpenMP/target_update_count_expression_codegen.c
new file mode 100644
index 0000000000000..df7811d078313
--- /dev/null
+++ b/clang/test/OpenMP/target_update_count_expression_codegen.c
@@ -0,0 +1,41 @@
+// 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 OMPIRBuilder.cpp correctly stores dimension count (not byte size)
+// in offload_sizes for NON_CONTIG entries with count expressions.
+//
+// This test verifies the IR change: offload_sizes contains [i64 2] for
+// NON_CONTIG entries (dimension count) and [i64 20] for non-NON_CONTIG entries
+// (byte size: 5 elements * 4 bytes).
+
+// CK-64-DAG: [[SIZE_NON_CONTIG:@.+]] = {{.+}}constant [1 x i64] [i64 2]
+// CK-32-DAG: [[SIZE_NON_CONTIG:@.+]] = {{.+}}constant [1 x i64] [i64 2]
+// For non-NON_CONTIG entries, offload_sizes contains byte size (5 elements * 4 bytes = 20)
+// CK-64-DAG: [[SIZE_CONTIG:@.+]] = {{.+}}constant [1 x i64] [i64 20]
+// CK-32-DAG: [[SIZE_CONTIG:@.+]] = {{.+}}constant [1 x i64] [i64 20]
+
+void test_non_contig_dimension_count() {
+ const int len = 10;
+ int data[10];
+ // CK-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr {{%.+}}, ptr {{%.+}}, ptr [[SIZE_NON_CONTIG]], ptr {{.+}}, ptr null, ptr null)
+ #pragma omp target update to(data[0:len/2:2])
+}
+
+void test_contig_byte_size() {
+ int data[10];
+ // CK-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr {{%.+}}, ptr {{%.+}}, ptr [[SIZE_CONTIG]], ptr {{.+}}, 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