[clang] [llvm] [Clang][OpenMP] Support expression semantics in `target update` fields with non-contiguous array sections (PR #175505)
Amit Tiwari via cfe-commits
cfe-commits at lists.llvm.org
Sun Jan 18 23:55:56 PST 2026
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/175505
>From 3ce6d1973f06635bcbf810e0b1909d60918bd390 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 716f8582dd7b2..fe6c755b0e504 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9365,16 +9365,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 0ffe45560dc0e9206486eafa7aeb8a9a11755844 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Thu, 4 Dec 2025 05:10:55 -0500
Subject: [PATCH 2/5] testcases_count_expression
---
.../strided_update_count_expression_from.c | 54 ++++++++
.../strided_update_count_expression_to.c | 72 ++++++++++
...xpression_stride_greater_than_count_from.c | 42 ++++++
...ontiguous_count_expression_variable_from.c | 123 +++++++++++++++++
...-contiguous_count_expression_variable_to.c | 125 ++++++++++++++++++
...n-contiguous_count_expression_zero_count.c | 43 ++++++
.../target_update_ptr_count_expression_from.c | 74 +++++++++++
.../target_update_ptr_count_expression_to.c | 82 ++++++++++++
...ate_strided_struct_count_expression_from.c | 86 ++++++++++++
...pdate_strided_struct_count_expression_to.c | 98 ++++++++++++++
10 files changed, 799 insertions(+)
create mode 100644 offload/test/offloading/strided_update_count_expression_from.c
create mode 100644 offload/test/offloading/strided_update_count_expression_to.c
create mode 100644 offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c
create mode 100644 offload/test/offloading/target_non-contiguous_count_expression_variable_from.c
create mode 100644 offload/test/offloading/target_non-contiguous_count_expression_variable_to.c
create mode 100644 offload/test/offloading/target_non-contiguous_count_expression_zero_count.c
create mode 100644 offload/test/offloading/target_update_ptr_count_expression_from.c
create mode 100644 offload/test/offloading/target_update_ptr_count_expression_to.c
create mode 100644 offload/test/offloading/target_update_strided_struct_count_expression_from.c
create mode 100644 offload/test/offloading/target_update_strided_struct_count_expression_to.c
diff --git a/offload/test/offloading/strided_update_count_expression_from.c b/offload/test/offloading/strided_update_count_expression_from.c
new file mode 100644
index 0000000000000..d33ba9e428af5
--- /dev/null
+++ b/offload/test/offloading/strided_update_count_expression_from.c
@@ -0,0 +1,54 @@
+// This test checks that "update from" clause in OpenMP is supported when the
+// elements are updated in a non-contiguous manner. This test checks that
+// #pragma omp target update from(data[0:len/2:2]) correctly updates only every
+// other element (stride 2) from the device to the host
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+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;
+ }
+ }
+ // Initial values
+ printf("original host array values:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+#pragma omp target data map(to : len, data[0 : len])
+ {
+// Modify arrays on device
+#pragma omp target
+ for (int i = 0; i < len; i++) {
+ data[i] += i;
+ }
+
+#pragma omp target update from(data[0 : len/2 : 2])
+ }
+ // CHECK: 0.000000
+ // CHECK: 1.000000
+ // CHECK: 4.000000
+ // CHECK: 3.000000
+ // CHECK: 8.000000
+ // CHECK: 5.000000
+ // CHECK: 12.000000
+ // CHECK: 7.000000
+ // CHECK-NOT: 2.000000
+ // CHECK-NOT: 6.000000
+ // CHECK-NOT: 10.000000
+ // CHECK-NOT: 14.000000
+
+ printf("from target array results:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+ return 0;
+}
diff --git a/offload/test/offloading/strided_update_count_expression_to.c b/offload/test/offloading/strided_update_count_expression_to.c
new file mode 100644
index 0000000000000..0b1f179c467e3
--- /dev/null
+++ b/offload/test/offloading/strided_update_count_expression_to.c
@@ -0,0 +1,72 @@
+// This test checks that "update to" clause in OpenMP is supported when the
+// elements are updated in a non-contiguous manner. This test checks that
+// #pragma omp target update to(data[0:len/2:2]) correctly updates only every
+// other element (stride 2) from the host to the device
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int len = 8;
+ double data[len];
+
+ // Initialize on host
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ }
+
+ // Initial values
+ printf("original host array values:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+#pragma omp target data map(tofrom : len, data[0 : len])
+ {
+ // Initialize device to 20
+#pragma omp target
+ for (int i = 0; i < len; i++) {
+ data[i] = 20.0;
+ }
+
+ // Modify host for strided elements
+ data[0] = 10.0;
+ data[2] = 10.0;
+ data[4] = 10.0;
+ data[6] = 10.0;
+
+#pragma omp target update to(data[0 : len/2 : 2])
+
+ // Verify on device by adding 5
+#pragma omp target
+ for (int i = 0; i < len; i++) {
+ data[i] += 5.0;
+ }
+ }
+
+ // CHECK: 0.000000
+ // CHECK: 1.000000
+ // CHECK: 2.000000
+ // CHECK: 3.000000
+ // CHECK: 4.000000
+ // CHECK: 5.000000
+ // CHECK: 6.000000
+ // CHECK: 7.000000
+
+ printf("device array values after update to:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+ // CHECK: 15.000000
+ // CHECK: 25.000000
+ // CHECK: 15.000000
+ // CHECK: 25.000000
+ // CHECK: 15.000000
+ // CHECK: 25.000000
+ // CHECK: 15.000000
+ // CHECK: 25.000000
+
+ return 0;
+}
diff --git a/offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c b/offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c
new file mode 100644
index 0000000000000..2677fe1310760
--- /dev/null
+++ b/offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c
@@ -0,0 +1,42 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+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;
+ }
+ }
+
+#pragma omp target data map(to : len, data[0 : len])
+ {
+#pragma omp target
+ for (int i = 0; i < len; i++) {
+ data[i] += i;
+ }
+
+ int small_count = 2;
+#pragma omp target update from(data[0 : small_count : 10])
+ }
+
+ printf("from target array results:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+ return 0;
+}
+
+// CHECK: from target array results:
+// CHECK: 0.000000
+// CHECK: 1.000000
+// CHECK: 2.000000
+// CHECK: 3.000000
+// CHECK: 4.000000
+// CHECK: 5.000000
+// CHECK: 6.000000
+// CHECK: 7.000000
diff --git a/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c b/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c
new file mode 100644
index 0000000000000..74fed75049cd2
--- /dev/null
+++ b/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c
@@ -0,0 +1,123 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int len = 16;
+ double data[len];
+ double data1[len], data2[len];
+
+ // Initialize data, data1, data2 on device
+#pragma omp target map(tofrom : len, data[0 : len], data1[0 : len], data2[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ data1[i] = i;
+ data2[i] = i * 10;
+ }
+ }
+
+#pragma omp target data map(to : len, data[0 : len], data1[0 : len], data2[0 : len])
+ {
+ // Device modifies arrays:
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] += i;
+ data1[i] += i;
+ data2[i] += 100;
+ }
+ }
+
+ int count = 4;
+ // indices: {0, 2, 4, 6}
+#pragma omp target update from(data[0 : count : 2])
+
+ // indices: {0, 2, 4, 6, 8, 10, 12, 14}
+#pragma omp target update from(data[0 : len/2 : 2])
+
+ // indices: {2, 4, 6, 8, 10, 12, 14}
+#pragma omp target update from(data[2 : len-4 : 2])
+
+int partial_count = 4;
+ // indices: {0, 3, 6, 9}
+#pragma omp target update from(data[0 : partial_count : 3])
+
+int count1 = 3;
+int count2 = 2;
+ // data1 indices: {0, 4, 8}
+ // data2 indices: {0, 5}
+#pragma omp target update from(data1[0 : count1 : 4], data2[0 : count2 : 5])
+ }
+
+ // Print results
+ printf("from target array results (data):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+ printf("from target array results (data1, data2):\n");
+ printf("data1:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data1[i]);
+ printf("data2:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data2[i]);
+ printf("\n");
+
+ return 0;
+}
+
+// CHECK: from target array results (data):
+// CHECK: 0.000000
+// CHECK: 1.000000
+// CHECK: 4.000000
+// CHECK: 6.000000
+// CHECK: 8.000000
+// CHECK: 5.000000
+// CHECK: 12.000000
+// CHECK: 7.000000
+// CHECK: 16.000000
+// CHECK: 18.000000
+// CHECK: 20.000000
+// CHECK: 11.000000
+// CHECK: 24.000000
+// CHECK: 13.000000
+// CHECK: 28.000000
+// CHECK: 15.000000
+
+// CHECK: from target array results (data1, data2):
+// CHECK: data1:
+// CHECK: 0.000000
+// CHECK: 1.000000
+// CHECK: 2.000000
+// CHECK: 3.000000
+// CHECK: 8.000000
+// CHECK: 5.000000
+// CHECK: 6.000000
+// CHECK: 7.000000
+// CHECK: 16.000000
+// CHECK: 9.000000
+// CHECK: 10.000000
+// CHECK: 11.000000
+// CHECK: 12.000000
+// CHECK: 13.000000
+// CHECK: 14.000000
+// CHECK: 15.000000
+// CHECK: data2:
+// CHECK: 100.000000
+// CHECK: 10.000000
+// CHECK: 20.000000
+// CHECK: 30.000000
+// CHECK: 40.000000
+// CHECK: 150.000000
+// CHECK: 60.000000
+// CHECK: 70.000000
+// CHECK: 80.000000
+// CHECK: 90.000000
+// CHECK: 100.000000
+// CHECK: 110.000000
+// CHECK: 120.000000
+// CHECK: 130.000000
+// CHECK: 140.000000
+// CHECK: 150.000000
diff --git a/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c b/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c
new file mode 100644
index 0000000000000..3cf53f0c206da
--- /dev/null
+++ b/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c
@@ -0,0 +1,125 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int len = 16;
+ double data[len];
+ double data1[len], data2[len];
+
+ // Initialize data, data1, data2 on device (and copy back to host)
+#pragma omp target map(tofrom : len, data[0 : len], data1[0 : len], data2[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ data1[i] = i;
+ data2[i] = i * 10;
+ }
+ }
+
+#pragma omp target data map(to : len, data[0 : len], data1[0 : len], data2[0 : len])
+ {
+ int count = 4;
+ // indices: {0, 2, 4, 6}
+#pragma omp target update to(data[0 : count : 2])
+
+ int half_len = len / 2;
+ // indices: {0, 2, 4, 6, 8, 10, 12, 14}
+#pragma omp target update to(data[0 : half_len : 2])
+
+ int dyn_count = len - 4;
+ // indices: {2, 4, 6, 8, 10, 12, 14}
+#pragma omp target update to(data[2 : dyn_count : 2])
+
+ int partial_count = 4;
+ // indices: {0, 3, 6, 9}
+#pragma omp target update to(data[0 : partial_count : 3])
+
+ int count1 = 3;
+ int count2 = 2;
+ // data1 indices: {0, 4, 8}
+ // data2 indices: {0, 5}
+#pragma omp target update to(data1[0 : count1 : 4], data2[0 : count2 : 5])
+
+ // Device modifies arrays
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] += i; // becomes 2*i on device
+ data1[i] += i; // becomes 2*i on device
+ data2[i] += 100; // becomes i*10 + 100 on device
+ }
+ }
+ }
+
+ // Print results
+ printf("from target array results (data):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+ printf("from target array results (data1, data2):\n");
+ printf("data1:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data1[i]);
+ printf("data2:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data2[i]);
+ printf("\n");
+
+ return 0;
+}
+
+// CHECK: from target array results (data):
+// CHECK: 0.000000
+// CHECK: 1.000000
+// CHECK: 2.000000
+// CHECK: 3.000000
+// CHECK: 4.000000
+// CHECK: 5.000000
+// CHECK: 6.000000
+// CHECK: 7.000000
+// CHECK: 8.000000
+// CHECK: 9.000000
+// CHECK: 10.000000
+// CHECK: 11.000000
+// CHECK: 12.000000
+// CHECK: 13.000000
+// CHECK: 14.000000
+// CHECK: 15.000000
+
+// CHECK: from target array results (data1, data2):
+// CHECK: data1:
+// CHECK: 0.000000
+// CHECK: 1.000000
+// CHECK: 2.000000
+// CHECK: 3.000000
+// CHECK: 4.000000
+// CHECK: 5.000000
+// CHECK: 6.000000
+// CHECK: 7.000000
+// CHECK: 8.000000
+// CHECK: 9.000000
+// CHECK: 10.000000
+// CHECK: 11.000000
+// CHECK: 12.000000
+// CHECK: 13.000000
+// CHECK: 14.000000
+// CHECK: 15.000000
+// CHECK: data2:
+// CHECK: 0.000000
+// CHECK: 10.000000
+// CHECK: 20.000000
+// CHECK: 30.000000
+// CHECK: 40.000000
+// CHECK: 50.000000
+// CHECK: 60.000000
+// CHECK: 70.000000
+// CHECK: 80.000000
+// CHECK: 90.000000
+// CHECK: 100.000000
+// CHECK: 110.000000
+// CHECK: 120.000000
+// CHECK: 130.000000
+// CHECK: 140.000000
+// CHECK: 150.000000
diff --git a/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c b/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c
new file mode 100644
index 0000000000000..1567e79814b84
--- /dev/null
+++ b/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c
@@ -0,0 +1,43 @@
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+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;
+ }
+ }
+
+#pragma omp target data map(to : len, data[0 : len])
+ {
+#pragma omp target
+ for (int i = 0; i < len; i++) {
+ data[i] += i;
+ }
+
+ int zero_count = 0;
+#pragma omp target update from(data[0 : zero_count : 2])
+ }
+
+ printf("from target array results:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+ return 0;
+}
+
+// CHECK: from target array results:
+// CHECK: 0.000000
+// CHECK: 1.000000
+// CHECK: 2.000000
+// CHECK: 3.000000
+// CHECK: 4.000000
+// CHECK: 5.000000
+// CHECK: 6.000000
+// CHECK: 7.000000
diff --git a/offload/test/offloading/target_update_ptr_count_expression_from.c b/offload/test/offloading/target_update_ptr_count_expression_from.c
new file mode 100644
index 0000000000000..2c5740b906f6c
--- /dev/null
+++ b/offload/test/offloading/target_update_ptr_count_expression_from.c
@@ -0,0 +1,74 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that "update from" clause in OpenMP supports strided
+// sections. #pragma omp target update from(result[0:len/2:2]) updates every other
+// element from device
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 32
+
+int main() {
+ double *result = (double *)calloc(N, sizeof(double));
+ int len = N;
+
+ printf("initial host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", result[i]);
+ printf("\n");
+
+#pragma omp target data map(to : result[0 : N])
+ {
+#pragma omp target map(alloc : result[0 : N])
+ for (int i = 0; i < N; i++)
+ result[i] += i;
+
+ // Update strided elements from device: even indices 0,2,4,...,30
+#pragma omp target update from(result[0 : len/2 : 2])
+ }
+
+ printf("after target update from (even indices up to 30 updated):\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", result[i]);
+ printf("\n");
+
+ // Expected: even indices i, odd indices 0
+ // CHECK: 0.000000
+ // CHECK: 0.000000
+ // CHECK: 2.000000
+ // CHECK: 0.000000
+ // CHECK: 4.000000
+ // CHECK: 0.000000
+ // CHECK: 6.000000
+ // CHECK: 0.000000
+ // CHECK: 8.000000
+ // CHECK: 0.000000
+ // CHECK: 10.000000
+ // CHECK: 0.000000
+ // CHECK: 12.000000
+ // CHECK: 0.000000
+ // CHECK: 14.000000
+ // CHECK: 0.000000
+ // CHECK: 16.000000
+ // CHECK: 0.000000
+ // CHECK: 18.000000
+ // CHECK: 0.000000
+ // CHECK: 20.000000
+ // CHECK: 0.000000
+ // CHECK: 22.000000
+ // CHECK: 0.000000
+ // CHECK: 24.000000
+ // CHECK: 0.000000
+ // CHECK: 26.000000
+ // CHECK: 0.000000
+ // CHECK: 28.000000
+ // CHECK: 0.000000
+ // CHECK: 30.000000
+ // CHECK: 0.000000
+ // CHECK-NOT: 1.000000
+ // CHECK-NOT: 3.000000
+ // CHECK-NOT: 31.000000
+
+ free(result);
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_ptr_count_expression_to.c b/offload/test/offloading/target_update_ptr_count_expression_to.c
new file mode 100644
index 0000000000000..3900592f10f93
--- /dev/null
+++ b/offload/test/offloading/target_update_ptr_count_expression_to.c
@@ -0,0 +1,82 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that "update to" clause in OpenMP supports strided sections.
+// #pragma omp target update to(result[0:len/2:2]) updates every other element
+// (stride 2)
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16
+
+int main() {
+ double *result = (double *)calloc(N, sizeof(double));
+ int len = N;
+
+ // Initialize on host
+ for (int i = 0; i < N; i++) {
+ result[i] = i;
+ }
+
+ // Initial values
+ printf("original host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", result[i]);
+ printf("\n");
+
+#pragma omp target data map(tofrom : result[0 : N])
+ {
+// Update strided elements to device: indices 0,2,4,6
+#pragma omp target update to(result[0 : len/2 : 2])
+
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++) {
+ result[i] += i;
+ }
+ }
+ }
+
+ printf("from target array results:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", result[i]);
+
+ // 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-NEXT: 10.000000
+ // CHECK-NEXT: 11.000000
+ // CHECK-NEXT: 12.000000
+ // CHECK-NEXT: 13.000000
+ // CHECK-NEXT: 14.000000
+ // CHECK-NEXT: 15.000000
+
+ // CHECK: from target array results:
+ // CHECK-NEXT: 0.000000
+ // CHECK-NEXT: 2.000000
+ // CHECK-NEXT: 4.000000
+ // CHECK-NEXT: 6.000000
+ // CHECK-NEXT: 8.000000
+ // CHECK-NEXT: 10.000000
+ // CHECK-NEXT: 12.000000
+ // CHECK-NEXT: 14.000000
+ // CHECK-NEXT: 16.000000
+ // CHECK-NEXT: 18.000000
+ // CHECK-NEXT: 20.000000
+ // CHECK-NEXT: 22.000000
+ // CHECK-NEXT: 24.000000
+ // CHECK-NEXT: 26.000000
+ // CHECK-NEXT: 28.000000
+ // CHECK-NEXT: 30.000000
+
+ free(result);
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_strided_struct_count_expression_from.c b/offload/test/offloading/target_update_strided_struct_count_expression_from.c
new file mode 100644
index 0000000000000..bb3e2ba577f7f
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_count_expression_from.c
@@ -0,0 +1,86 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that "update from" with user-defined mapper supports strided
+// sections using fixed-size arrays in structs.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16
+
+typedef struct {
+ double data[N];
+ size_t 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;
+
+ 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("%f\n", s.data[i]);
+ printf("\n");
+
+#pragma omp target data map(mapper(custom), tofrom : s)
+ {
+// Execute on device with explicit mapper
+#pragma omp target map(mapper(custom), tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += i;
+ }
+ }
+
+// Update strided elements from device: indices 0,2,4,6,8,10,12,14
+#pragma omp target update from(s.data[0 : s.len/2 : 2])
+ }
+
+ printf("from target array results:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", s.data[i]);
+
+ // 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-NEXT: 10.000000
+ // CHECK-NEXT: 11.000000
+ // CHECK-NEXT: 12.000000
+ // CHECK-NEXT: 13.000000
+ // CHECK-NEXT: 14.000000
+ // CHECK-NEXT: 15.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-NEXT: 20.000000
+ // CHECK-NEXT: 11.000000
+ // CHECK-NEXT: 24.000000
+ // CHECK-NEXT: 13.000000
+ // CHECK-NEXT: 28.000000
+ // CHECK-NEXT: 15.000000
+
+ return 0;
+}
diff --git a/offload/test/offloading/target_update_strided_struct_count_expression_to.c b/offload/test/offloading/target_update_strided_struct_count_expression_to.c
new file mode 100644
index 0000000000000..c4bc7faf367a0
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_count_expression_to.c
@@ -0,0 +1,98 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that "update to" with struct member arrays supports strided
+// sections using fixed-size arrays in structs.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16
+
+typedef struct {
+ double data[N];
+ int len;
+} T;
+
+int main() {
+ T s;
+ s.len = N;
+
+ // Initialize struct array on host with simple sequential values
+ for (int i = 0; i < N; i++) {
+ s.data[i] = i;
+ }
+
+ printf("original host struct 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)
+ {
+// Initialize device struct array to 20
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = 20.0;
+ }
+ }
+
+ // Modify host struct data for strided elements (set to 10)
+ for (int i = 0; i < 8; i++) {
+ s.data[i * 2] = 10.0; // Set even indices to 10
+ }
+
+// indices 0,2,4,6,8,10,12,14
+#pragma omp target update to(s.data[0 : s.len/2 : 2])
+
+// Execute on device - add 5 to verify update worked
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += 5.0;
+ }
+ }
+ }
+
+ printf("after target update to struct:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f\n", s.data[i]);
+
+ // CHECK: original host struct 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: after target update to struct:
+ // 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
+
+ return 0;
+}
>From bbb8410ab124efb1c2fedbb5a707b8eb609f1190 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Thu, 4 Dec 2025 06:34:09 -0500
Subject: [PATCH 3/5] 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 b8ee701c482bb..32b76f5d74d46 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7907,12 +7907,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 25ad443780e13987693474fda091713c83844988 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 12 Jan 2026 03:12:42 -0500
Subject: [PATCH 4/5] cleanup
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 164 +++++-----
...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 +++++
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 4 +-
.../strided_update_count_expression.c | 133 ++++++++
.../strided_update_count_expression_complex.c | 289 +++++++++++++++++
.../strided_update_count_expression_from.c | 54 ----
.../strided_update_count_expression_misc.c | 99 ++++++
.../strided_update_count_expression_to.c | 72 -----
..._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 ++++++
...xpression_stride_greater_than_count_from.c | 42 ---
...ontiguous_count_expression_variable_from.c | 123 --------
...-contiguous_count_expression_variable_to.c | 125 --------
...n-contiguous_count_expression_zero_count.c | 43 ---
.../target_update_ptr_count_expression.c | 99 ++++++
.../target_update_ptr_count_expression_from.c | 74 -----
.../target_update_ptr_count_expression_to.c | 82 -----
...get_update_ptr_variable_count_and_stride.c | 94 ++++++
.../target_update_ptr_variable_stride.c | 95 ++++++
...t_update_strided_struct_count_expression.c | 97 ++++++
...ate_strided_struct_count_expression_from.c | 86 -----
...pdate_strided_struct_count_expression_to.c | 98 ------
...strided_struct_variable_count_and_stride.c | 96 ++++++
...et_update_strided_struct_variable_stride.c | 95 ++++++
32 files changed, 2482 insertions(+), 886 deletions(-)
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
delete mode 100644 offload/test/offloading/strided_update_count_expression_from.c
create mode 100644 offload/test/offloading/strided_update_count_expression_misc.c
delete mode 100644 offload/test/offloading/strided_update_count_expression_to.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
delete mode 100644 offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c
delete mode 100644 offload/test/offloading/target_non-contiguous_count_expression_variable_from.c
delete mode 100644 offload/test/offloading/target_non-contiguous_count_expression_variable_to.c
delete mode 100644 offload/test/offloading/target_non-contiguous_count_expression_zero_count.c
create mode 100644 offload/test/offloading/target_update_ptr_count_expression.c
delete mode 100644 offload/test/offloading/target_update_ptr_count_expression_from.c
delete mode 100644 offload/test/offloading/target_update_ptr_count_expression_to.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
delete mode 100644 offload/test/offloading/target_update_strided_struct_count_expression_from.c
delete mode 100644 offload/test/offloading/target_update_strided_struct_count_expression_to.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/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 32b76f5d74d46..bd7189dbf23a4 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -492,11 +492,11 @@ enum OpenMPLocationFlags : unsigned {
/// member */
/// kmp_int32 reserved_2; /**< not really used in Fortran any more;
/// see above */
-///#if USE_ITT_BUILD
+/// #if USE_ITT_BUILD
/// /* but currently used for storing
/// region-specific ITT */
/// /* contextual information. */
-///#endif /* USE_ITT_BUILD */
+/// #endif /* USE_ITT_BUILD */
/// kmp_int32 reserved_3; /**< source[4] in Fortran, do not use for
/// C++ */
/// char const *psource; /**< String describing the source location.
@@ -716,16 +716,16 @@ static void EmitOMPAggregateInit(CodeGenFunction &CGF, Address DestAddr,
if (DRD) {
// Shift the address forward by one element.
- llvm::Value *SrcElementNext = CGF.Builder.CreateConstGEP1_32(
- SrcAddr.getElementType(), SrcElementPHI, /*Idx0=*/1,
- "omp.arraycpy.dest.element");
+ llvm::Value *SrcElementNext =
+ CGF.Builder.CreateConstGEP1_32(SrcAddr.getElementType(), SrcElementPHI,
+ /*Idx0=*/1, "omp.arraycpy.dest.element");
SrcElementPHI->addIncoming(SrcElementNext, CGF.Builder.GetInsertBlock());
}
// Shift the address forward by one element.
- llvm::Value *DestElementNext = CGF.Builder.CreateConstGEP1_32(
- DestAddr.getElementType(), DestElementPHI, /*Idx0=*/1,
- "omp.arraycpy.dest.element");
+ llvm::Value *DestElementNext =
+ CGF.Builder.CreateConstGEP1_32(DestAddr.getElementType(), DestElementPHI,
+ /*Idx0=*/1, "omp.arraycpy.dest.element");
// Check whether we've reached the end.
llvm::Value *Done =
CGF.Builder.CreateICmpEQ(DestElementNext, DestEnd, "omp.arraycpy.done");
@@ -975,8 +975,8 @@ Address ReductionCodeGen::adjustPrivateAddress(CodeGenFunction &CGF, unsigned N,
llvm::Value *PrivatePointer =
CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
PrivateAddr.emitRawPointer(CGF), SharedAddr.getType());
- llvm::Value *Ptr = CGF.Builder.CreateGEP(
- SharedAddr.getElementType(), PrivatePointer, Adjustment);
+ llvm::Value *Ptr = CGF.Builder.CreateGEP(SharedAddr.getElementType(),
+ PrivatePointer, Adjustment);
return castToBase(CGF, OrigVD->getType(),
SharedAddresses[N].first.getType(),
OriginalBaseLValue.getAddress(), Ptr);
@@ -1605,12 +1605,11 @@ Address CGOpenMPRuntime::getAddrOfThreadPrivate(CodeGenFunction &CGF,
CGF.Builder.CreatePointerCast(VDAddr.emitRawPointer(CGF), CGM.Int8PtrTy),
CGM.getSize(CGM.GetTargetTypeStoreSize(VarTy)),
getOrCreateThreadPrivateCache(VD)};
- return Address(
- CGF.EmitRuntimeCall(
- OMPBuilder.getOrCreateRuntimeFunction(
- CGM.getModule(), OMPRTL___kmpc_threadprivate_cached),
- Args),
- CGF.Int8Ty, VDAddr.getAlignment());
+ return Address(CGF.EmitRuntimeCall(
+ OMPBuilder.getOrCreateRuntimeFunction(
+ CGM.getModule(), OMPRTL___kmpc_threadprivate_cached),
+ Args),
+ CGF.Int8Ty, VDAddr.getAlignment());
}
void CGOpenMPRuntime::emitThreadPrivateVarInit(
@@ -1635,8 +1634,8 @@ void CGOpenMPRuntime::emitThreadPrivateVarInit(
}
llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition(
- const VarDecl *VD, Address VDAddr, SourceLocation Loc,
- bool PerformInit, CodeGenFunction *CGF) {
+ const VarDecl *VD, Address VDAddr, SourceLocation Loc, bool PerformInit,
+ CodeGenFunction *CGF) {
if (CGM.getLangOpts().OpenMPUseTLS &&
CGM.getContext().getTargetInfo().isTLSSupported())
return nullptr;
@@ -1698,7 +1697,8 @@ llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition(
auto NL = ApplyDebugLocation::CreateEmpty(DtorCGF);
DtorCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, Fn, FI, Args,
Loc, Loc);
- // Create a scope with an artificial location for the body of this function.
+ // Create a scope with an artificial location for the body of this
+ // function.
auto AL = ApplyDebugLocation::CreateArtificial(DtorCGF);
llvm::Value *ArgVal = DtorCGF.EmitLoadOfScalar(
DtorCGF.GetAddrOfLocalVar(&Dst),
@@ -1942,8 +1942,7 @@ Address CGOpenMPRuntime::emitThreadIDAddress(CodeGenFunction &CGF,
QualType Int32Ty =
CGF.getContext().getIntTypeForBitwidth(/*DestWidth*/ 32, /*Signed*/ true);
Address ThreadIDTemp = CGF.CreateMemTemp(Int32Ty, /*Name*/ ".threadid_temp.");
- CGF.EmitStoreOfScalar(ThreadID,
- CGF.MakeAddrLValue(ThreadIDTemp, Int32Ty));
+ CGF.EmitStoreOfScalar(ThreadID, CGF.MakeAddrLValue(ThreadIDTemp, Int32Ty));
return ThreadIDTemp;
}
@@ -2450,8 +2449,8 @@ bool CGOpenMPRuntime::isStaticChunked(OpenMPScheduleClauseKind ScheduleKind,
return Schedule == OMP_sch_static_chunked;
}
-bool CGOpenMPRuntime::isStaticChunked(
- OpenMPDistScheduleClauseKind ScheduleKind, bool Chunked) const {
+bool CGOpenMPRuntime::isStaticChunked(OpenMPDistScheduleClauseKind ScheduleKind,
+ bool Chunked) const {
OpenMPSchedType Schedule = getRuntimeSchedule(ScheduleKind, Chunked);
return Schedule == OMP_dist_sch_static_chunked;
}
@@ -2615,10 +2614,10 @@ void CGOpenMPRuntime::emitForStaticInit(CodeGenFunction &CGF,
ScheduleKind.Schedule, Values.Chunk != nullptr, Values.Ordered);
assert((isOpenMPWorksharingDirective(DKind) || (DKind == OMPD_loop)) &&
"Expected loop-based or sections-based directive.");
- llvm::Value *UpdatedLocation = emitUpdateLocation(CGF, Loc,
- isOpenMPLoopDirective(DKind)
- ? OMP_IDENT_WORK_LOOP
- : OMP_IDENT_WORK_SECTIONS);
+ llvm::Value *UpdatedLocation = emitUpdateLocation(
+ CGF, Loc,
+ isOpenMPLoopDirective(DKind) ? OMP_IDENT_WORK_LOOP
+ : OMP_IDENT_WORK_SECTIONS);
llvm::Value *ThreadId = getThreadID(CGF, Loc);
llvm::FunctionCallee StaticInitFunction =
OMPBuilder.createForStaticInitFunction(Values.IVSize, Values.IVSigned,
@@ -2693,9 +2692,8 @@ void CGOpenMPRuntime::emitForOrderedIterationEnd(CodeGenFunction &CGF,
llvm::Value *CGOpenMPRuntime::emitForNext(CodeGenFunction &CGF,
SourceLocation Loc, unsigned IVSize,
- bool IVSigned, Address IL,
- Address LB, Address UB,
- Address ST) {
+ bool IVSigned, Address IL, Address LB,
+ Address UB, Address ST) {
// Call __kmpc_dispatch_next(
// ident_t *loc, kmp_int32 tid, kmp_int32 *p_lastiter,
// kmp_int[32|64] *p_lower, kmp_int[32|64] *p_upper,
@@ -2891,8 +2889,8 @@ static bool isAllocatableDecl(const VarDecl *VD) {
!AA->getAllocator());
}
-static RecordDecl *
-createPrivatesRecordDecl(CodeGenModule &CGM, ArrayRef<PrivateDataTy> Privates) {
+static RecordDecl *createPrivatesRecordDecl(CodeGenModule &CGM,
+ ArrayRef<PrivateDataTy> Privates) {
if (!Privates.empty()) {
ASTContext &C = CGM.getContext();
// Build struct .kmp_privates_t. {
@@ -3396,7 +3394,6 @@ static bool checkInitIsRequired(CodeGenFunction &CGF,
return InitRequired;
}
-
/// Emit task_dup function (for initialization of
/// private/firstprivate/lastprivate vars and last_iter flag)
/// \code
@@ -3760,10 +3757,14 @@ CGOpenMPRuntime::emitTaskInit(CodeGenFunction &CGF, SourceLocation Loc,
: CGF.Builder.getInt32(Data.Final.getInt() ? FinalFlag : 0);
TaskFlags = CGF.Builder.CreateOr(TaskFlags, CGF.Builder.getInt32(Flags));
llvm::Value *SharedsSize = CGM.getSize(C.getTypeSizeInChars(SharedsTy));
- SmallVector<llvm::Value *, 8> AllocArgs = {emitUpdateLocation(CGF, Loc),
- getThreadID(CGF, Loc), TaskFlags, KmpTaskTWithPrivatesTySize,
- SharedsSize, CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
- TaskEntry, KmpRoutineEntryPtrTy)};
+ SmallVector<llvm::Value *, 8> AllocArgs = {
+ emitUpdateLocation(CGF, Loc),
+ getThreadID(CGF, Loc),
+ TaskFlags,
+ KmpTaskTWithPrivatesTySize,
+ SharedsSize,
+ CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(TaskEntry,
+ KmpRoutineEntryPtrTy)};
llvm::Value *NewTask;
if (D.hasClausesOfKind<OMPNowaitClause>()) {
// Check if we have any device clause associated with the directive.
@@ -3954,13 +3955,13 @@ CGOpenMPRuntime::emitTaskInit(CodeGenFunction &CGF, SourceLocation Loc,
// Copy shareds if there are any.
Address KmpTaskSharedsPtr = Address::invalid();
if (!SharedsTy->castAsRecordDecl()->field_empty()) {
- KmpTaskSharedsPtr = Address(
- CGF.EmitLoadOfScalar(
- CGF.EmitLValueForField(
- TDBase,
- *std::next(KmpTaskTQTyRD->field_begin(), KmpTaskTShareds)),
- Loc),
- CGF.Int8Ty, CGM.getNaturalTypeAlignment(SharedsTy));
+ KmpTaskSharedsPtr =
+ Address(CGF.EmitLoadOfScalar(
+ CGF.EmitLValueForField(
+ TDBase, *std::next(KmpTaskTQTyRD->field_begin(),
+ KmpTaskTShareds)),
+ Loc),
+ CGF.Int8Ty, CGM.getNaturalTypeAlignment(SharedsTy));
LValue Dest = CGF.MakeAddrLValue(KmpTaskSharedsPtr, SharedsTy);
LValue Src = CGF.MakeAddrLValue(Shareds, SharedsTy);
CGF.EmitAggregateCopy(Dest, Src, SharedsTy, AggValueSlot::DoesNotOverlap);
@@ -4543,7 +4544,7 @@ void CGOpenMPRuntime::emitTaskCall(CodeGenFunction &CGF, SourceLocation Loc,
// list is not empty
llvm::Value *ThreadID = getThreadID(CGF, Loc);
llvm::Value *UpLoc = emitUpdateLocation(CGF, Loc);
- llvm::Value *TaskArgs[] = { UpLoc, ThreadID, NewTask };
+ llvm::Value *TaskArgs[] = {UpLoc, ThreadID, NewTask};
llvm::Value *DepTaskArgs[7];
if (!Data.Dependences.empty()) {
DepTaskArgs[0] = UpLoc;
@@ -4788,12 +4789,12 @@ static void EmitOMPAggregateReduction(
Scope.ForceCleanup();
// Shift the address forward by one element.
- llvm::Value *LHSElementNext = CGF.Builder.CreateConstGEP1_32(
- LHSAddr.getElementType(), LHSElementPHI, /*Idx0=*/1,
- "omp.arraycpy.dest.element");
- llvm::Value *RHSElementNext = CGF.Builder.CreateConstGEP1_32(
- RHSAddr.getElementType(), RHSElementPHI, /*Idx0=*/1,
- "omp.arraycpy.src.element");
+ llvm::Value *LHSElementNext =
+ CGF.Builder.CreateConstGEP1_32(LHSAddr.getElementType(), LHSElementPHI,
+ /*Idx0=*/1, "omp.arraycpy.dest.element");
+ llvm::Value *RHSElementNext =
+ CGF.Builder.CreateConstGEP1_32(RHSAddr.getElementType(), RHSElementPHI,
+ /*Idx0=*/1, "omp.arraycpy.src.element");
// Check whether we've reached the end.
llvm::Value *Done =
CGF.Builder.CreateICmpEQ(LHSElementNext, LHSEnd, "omp.arraycpy.done");
@@ -5743,7 +5744,7 @@ llvm::Value *CGOpenMPRuntime::emitTaskReductionInit(
const FieldDecl *SharedFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *OrigFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *SizeFD = addFieldToRecordDecl(C, RD, C.getSizeType());
- const FieldDecl *InitFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
+ const FieldDecl *InitFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *FiniFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *CombFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *FlagsFD = addFieldToRecordDecl(
@@ -6253,7 +6254,7 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
/// Checks if the expression is constant or does not have non-trivial function
/// calls.
-static bool isTrivial(ASTContext &Ctx, const Expr * E) {
+static bool isTrivial(ASTContext &Ctx, const Expr *E) {
// We can skip constant expressions.
// We can skip expressions with trivial calls or simple expressions.
return (E->isEvaluatable(Ctx, Expr::SE_AllowUndefinedBehavior) ||
@@ -6448,10 +6449,11 @@ llvm::Value *CGOpenMPRuntime::emitNumTeamsForTargetDirective(
const auto *CS = D.getInnermostCapturedStmt();
CGOpenMPInnerExprInfo CGInfo(CGF, *CS);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
- llvm::Value *NumTeamsVal = CGF.EmitScalarExpr(NumTeams,
- /*IgnoreResultAssign*/ true);
+ llvm::Value *NumTeamsVal =
+ CGF.EmitScalarExpr(NumTeams,
+ /*IgnoreResultAssign*/ true);
return Bld.CreateIntCast(NumTeamsVal, CGF.Int32Ty,
- /*isSigned=*/true);
+ /*isSigned=*/true);
}
case OMPD_target_teams:
case OMPD_target_teams_distribute:
@@ -6459,10 +6461,11 @@ llvm::Value *CGOpenMPRuntime::emitNumTeamsForTargetDirective(
case OMPD_target_teams_distribute_parallel_for:
case OMPD_target_teams_distribute_parallel_for_simd: {
CodeGenFunction::RunCleanupsScope NumTeamsScope(CGF);
- llvm::Value *NumTeamsVal = CGF.EmitScalarExpr(NumTeams,
- /*IgnoreResultAssign*/ true);
+ llvm::Value *NumTeamsVal =
+ CGF.EmitScalarExpr(NumTeams,
+ /*IgnoreResultAssign*/ true);
return Bld.CreateIntCast(NumTeamsVal, CGF.Int32Ty,
- /*isSigned=*/true);
+ /*isSigned=*/true);
}
default:
break;
@@ -7907,9 +7910,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();
@@ -8129,8 +8132,8 @@ class MappableExprsHandler {
LowestElem, CGF.VoidPtrTy, CGF.Int8Ty),
TypeSize.getQuantity() - 1);
PartialStruct.HighestElem = {
- std::numeric_limits<decltype(
- PartialStruct.HighestElem.first)>::max(),
+ std::numeric_limits<
+ decltype(PartialStruct.HighestElem.first)>::max(),
HB};
PartialStruct.Base = BP;
PartialStruct.LB = LB;
@@ -11264,7 +11267,7 @@ bool CGOpenMPRuntime::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
return false;
const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
- switch(A->getAllocatorType()) {
+ switch (A->getAllocatorType()) {
case OMPAllocateDeclAttr::OMPNullMemAlloc:
case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
// Not supported, fallback to the default mem space.
@@ -11368,7 +11371,8 @@ void CGOpenMPRuntime::emitNumTeamsClause(CodeGenFunction &CGF,
CGF.CGM.Int32Ty, /* isSigned = */ true)
: CGF.Builder.getInt32(0);
- // Build call __kmpc_push_num_teamss(&loc, global_tid, num_teams, thread_limit)
+ // Build call __kmpc_push_num_teamss(&loc, global_tid, num_teams,
+ // thread_limit)
llvm::Value *PushNumTeamsArgs[] = {RTLoc, getThreadID(CGF, Loc), NumTeamsVal,
ThreadLimitVal};
CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
@@ -11677,7 +11681,7 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
}
namespace {
- /// Kind of parameter in a function with 'declare simd' directive.
+/// Kind of parameter in a function with 'declare simd' directive.
enum ParamKindTy {
Linear,
LinearRef,
@@ -11803,18 +11807,10 @@ emitX86DeclareSimdFunction(const FunctionDecl *FD, llvm::Function *Fn,
unsigned VecRegSize;
};
ISADataTy ISAData[] = {
- {
- 'b', 128
- }, // SSE
- {
- 'c', 256
- }, // AVX
- {
- 'd', 256
- }, // AVX2
- {
- 'e', 512
- }, // AVX512
+ {'b', 128}, // SSE
+ {'c', 256}, // AVX
+ {'d', 256}, // AVX2
+ {'e', 512}, // AVX512
};
llvm::SmallVector<char, 2> Masked;
switch (State) {
@@ -12795,7 +12791,8 @@ Address CGOpenMPRuntime::emitLastprivateConditionalInit(CodeGenFunction &CGF,
FiredField = addFieldToRecordDecl(C, RD, C.CharTy);
RD->completeDefinition();
NewType = C.getCanonicalTagType(RD);
- Address Addr = CGF.CreateMemTemp(NewType, C.getDeclAlign(VD), VD->getName());
+ Address Addr =
+ CGF.CreateMemTemp(NewType, C.getDeclAlign(VD), VD->getName());
BaseLVal = CGF.MakeAddrLValue(Addr, NewType, AlignmentSource::Decl);
I->getSecond().try_emplace(VD, NewType, VDField, FiredField, BaseLVal);
} else {
@@ -12804,8 +12801,7 @@ Address CGOpenMPRuntime::emitLastprivateConditionalInit(CodeGenFunction &CGF,
FiredField = std::get<2>(VI->getSecond());
BaseLVal = std::get<3>(VI->getSecond());
}
- LValue FiredLVal =
- CGF.EmitLValueForField(BaseLVal, FiredField);
+ LValue FiredLVal = CGF.EmitLValueForField(BaseLVal, FiredField);
CGF.EmitStoreOfScalar(
llvm::ConstantInt::getNullValue(CGF.ConvertTypeForMem(C.CharTy)),
FiredLVal);
@@ -12992,7 +12988,7 @@ void CGOpenMPRuntime::checkAndEmitLastprivateConditional(CodeGenFunction &CGF,
assert(It != LastprivateConditionalToTypes[FoundFn].end() &&
"Lastprivate conditional is not found in outer region.");
QualType StructTy = std::get<0>(It->getSecond());
- const FieldDecl* FiredDecl = std::get<2>(It->getSecond());
+ const FieldDecl *FiredDecl = std::get<2>(It->getSecond());
LValue PrivLVal = CGF.EmitLValue(FoundE);
Address StructAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
PrivLVal.getAddress(),
@@ -13346,9 +13342,7 @@ bool CGOpenMPSIMDRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
llvm_unreachable("Not supported in SIMD-only mode");
}
-bool CGOpenMPSIMDRuntime::emitTargetGlobal(GlobalDecl GD) {
- return false;
-}
+bool CGOpenMPSIMDRuntime::emitTargetGlobal(GlobalDecl GD) { return false; }
void CGOpenMPSIMDRuntime::emitTeamsCall(CodeGenFunction &CGF,
const OMPExecutableDirective &D,
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/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index fe6c755b0e504..418c6142380eb 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9368,8 +9368,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..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_from.c b/offload/test/offloading/strided_update_count_expression_from.c
deleted file mode 100644
index d33ba9e428af5..0000000000000
--- a/offload/test/offloading/strided_update_count_expression_from.c
+++ /dev/null
@@ -1,54 +0,0 @@
-// This test checks that "update from" clause in OpenMP is supported when the
-// elements are updated in a non-contiguous manner. This test checks that
-// #pragma omp target update from(data[0:len/2:2]) correctly updates only every
-// other element (stride 2) from the device to the host
-
-// RUN: %libomptarget-compile-run-and-check-generic
-#include <omp.h>
-#include <stdio.h>
-
-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;
- }
- }
- // Initial values
- printf("original host array values:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data[i]);
- printf("\n");
-
-#pragma omp target data map(to : len, data[0 : len])
- {
-// Modify arrays on device
-#pragma omp target
- for (int i = 0; i < len; i++) {
- data[i] += i;
- }
-
-#pragma omp target update from(data[0 : len/2 : 2])
- }
- // CHECK: 0.000000
- // CHECK: 1.000000
- // CHECK: 4.000000
- // CHECK: 3.000000
- // CHECK: 8.000000
- // CHECK: 5.000000
- // CHECK: 12.000000
- // CHECK: 7.000000
- // CHECK-NOT: 2.000000
- // CHECK-NOT: 6.000000
- // CHECK-NOT: 10.000000
- // CHECK-NOT: 14.000000
-
- printf("from target array results:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data[i]);
- printf("\n");
-
- return 0;
-}
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_count_expression_to.c b/offload/test/offloading/strided_update_count_expression_to.c
deleted file mode 100644
index 0b1f179c467e3..0000000000000
--- a/offload/test/offloading/strided_update_count_expression_to.c
+++ /dev/null
@@ -1,72 +0,0 @@
-// This test checks that "update to" clause in OpenMP is supported when the
-// elements are updated in a non-contiguous manner. This test checks that
-// #pragma omp target update to(data[0:len/2:2]) correctly updates only every
-// other element (stride 2) from the host to the device
-
-// RUN: %libomptarget-compile-run-and-check-generic
-#include <omp.h>
-#include <stdio.h>
-
-int main() {
- int len = 8;
- double data[len];
-
- // Initialize on host
- for (int i = 0; i < len; i++) {
- data[i] = i;
- }
-
- // Initial values
- printf("original host array values:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data[i]);
- printf("\n");
-
-#pragma omp target data map(tofrom : len, data[0 : len])
- {
- // Initialize device to 20
-#pragma omp target
- for (int i = 0; i < len; i++) {
- data[i] = 20.0;
- }
-
- // Modify host for strided elements
- data[0] = 10.0;
- data[2] = 10.0;
- data[4] = 10.0;
- data[6] = 10.0;
-
-#pragma omp target update to(data[0 : len/2 : 2])
-
- // Verify on device by adding 5
-#pragma omp target
- for (int i = 0; i < len; i++) {
- data[i] += 5.0;
- }
- }
-
- // CHECK: 0.000000
- // CHECK: 1.000000
- // CHECK: 2.000000
- // CHECK: 3.000000
- // CHECK: 4.000000
- // CHECK: 5.000000
- // CHECK: 6.000000
- // CHECK: 7.000000
-
- printf("device array values after update to:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data[i]);
- printf("\n");
-
- // CHECK: 15.000000
- // CHECK: 25.000000
- // CHECK: 15.000000
- // CHECK: 25.000000
- // CHECK: 15.000000
- // CHECK: 25.000000
- // CHECK: 15.000000
- // CHECK: 25.000000
-
- 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..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_non-contiguous_count_expression_stride_greater_than_count_from.c b/offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c
deleted file mode 100644
index 2677fe1310760..0000000000000
--- a/offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c
+++ /dev/null
@@ -1,42 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-#include <omp.h>
-#include <stdio.h>
-
-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;
- }
- }
-
-#pragma omp target data map(to : len, data[0 : len])
- {
-#pragma omp target
- for (int i = 0; i < len; i++) {
- data[i] += i;
- }
-
- int small_count = 2;
-#pragma omp target update from(data[0 : small_count : 10])
- }
-
- printf("from target array results:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data[i]);
- printf("\n");
-
- return 0;
-}
-
-// CHECK: from target array results:
-// CHECK: 0.000000
-// CHECK: 1.000000
-// CHECK: 2.000000
-// CHECK: 3.000000
-// CHECK: 4.000000
-// CHECK: 5.000000
-// CHECK: 6.000000
-// CHECK: 7.000000
diff --git a/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c b/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c
deleted file mode 100644
index 74fed75049cd2..0000000000000
--- a/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c
+++ /dev/null
@@ -1,123 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-#include <omp.h>
-#include <stdio.h>
-
-int main() {
- int len = 16;
- double data[len];
- double data1[len], data2[len];
-
- // Initialize data, data1, data2 on device
-#pragma omp target map(tofrom : len, data[0 : len], data1[0 : len], data2[0 : len])
- {
- for (int i = 0; i < len; i++) {
- data[i] = i;
- data1[i] = i;
- data2[i] = i * 10;
- }
- }
-
-#pragma omp target data map(to : len, data[0 : len], data1[0 : len], data2[0 : len])
- {
- // Device modifies arrays:
-#pragma omp target
- {
- for (int i = 0; i < len; i++) {
- data[i] += i;
- data1[i] += i;
- data2[i] += 100;
- }
- }
-
- int count = 4;
- // indices: {0, 2, 4, 6}
-#pragma omp target update from(data[0 : count : 2])
-
- // indices: {0, 2, 4, 6, 8, 10, 12, 14}
-#pragma omp target update from(data[0 : len/2 : 2])
-
- // indices: {2, 4, 6, 8, 10, 12, 14}
-#pragma omp target update from(data[2 : len-4 : 2])
-
-int partial_count = 4;
- // indices: {0, 3, 6, 9}
-#pragma omp target update from(data[0 : partial_count : 3])
-
-int count1 = 3;
-int count2 = 2;
- // data1 indices: {0, 4, 8}
- // data2 indices: {0, 5}
-#pragma omp target update from(data1[0 : count1 : 4], data2[0 : count2 : 5])
- }
-
- // Print results
- printf("from target array results (data):\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data[i]);
- printf("\n");
-
- printf("from target array results (data1, data2):\n");
- printf("data1:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data1[i]);
- printf("data2:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data2[i]);
- printf("\n");
-
- return 0;
-}
-
-// CHECK: from target array results (data):
-// CHECK: 0.000000
-// CHECK: 1.000000
-// CHECK: 4.000000
-// CHECK: 6.000000
-// CHECK: 8.000000
-// CHECK: 5.000000
-// CHECK: 12.000000
-// CHECK: 7.000000
-// CHECK: 16.000000
-// CHECK: 18.000000
-// CHECK: 20.000000
-// CHECK: 11.000000
-// CHECK: 24.000000
-// CHECK: 13.000000
-// CHECK: 28.000000
-// CHECK: 15.000000
-
-// CHECK: from target array results (data1, data2):
-// CHECK: data1:
-// CHECK: 0.000000
-// CHECK: 1.000000
-// CHECK: 2.000000
-// CHECK: 3.000000
-// CHECK: 8.000000
-// CHECK: 5.000000
-// CHECK: 6.000000
-// CHECK: 7.000000
-// CHECK: 16.000000
-// CHECK: 9.000000
-// CHECK: 10.000000
-// CHECK: 11.000000
-// CHECK: 12.000000
-// CHECK: 13.000000
-// CHECK: 14.000000
-// CHECK: 15.000000
-// CHECK: data2:
-// CHECK: 100.000000
-// CHECK: 10.000000
-// CHECK: 20.000000
-// CHECK: 30.000000
-// CHECK: 40.000000
-// CHECK: 150.000000
-// CHECK: 60.000000
-// CHECK: 70.000000
-// CHECK: 80.000000
-// CHECK: 90.000000
-// CHECK: 100.000000
-// CHECK: 110.000000
-// CHECK: 120.000000
-// CHECK: 130.000000
-// CHECK: 140.000000
-// CHECK: 150.000000
diff --git a/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c b/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c
deleted file mode 100644
index 3cf53f0c206da..0000000000000
--- a/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c
+++ /dev/null
@@ -1,125 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-#include <omp.h>
-#include <stdio.h>
-
-int main() {
- int len = 16;
- double data[len];
- double data1[len], data2[len];
-
- // Initialize data, data1, data2 on device (and copy back to host)
-#pragma omp target map(tofrom : len, data[0 : len], data1[0 : len], data2[0 : len])
- {
- for (int i = 0; i < len; i++) {
- data[i] = i;
- data1[i] = i;
- data2[i] = i * 10;
- }
- }
-
-#pragma omp target data map(to : len, data[0 : len], data1[0 : len], data2[0 : len])
- {
- int count = 4;
- // indices: {0, 2, 4, 6}
-#pragma omp target update to(data[0 : count : 2])
-
- int half_len = len / 2;
- // indices: {0, 2, 4, 6, 8, 10, 12, 14}
-#pragma omp target update to(data[0 : half_len : 2])
-
- int dyn_count = len - 4;
- // indices: {2, 4, 6, 8, 10, 12, 14}
-#pragma omp target update to(data[2 : dyn_count : 2])
-
- int partial_count = 4;
- // indices: {0, 3, 6, 9}
-#pragma omp target update to(data[0 : partial_count : 3])
-
- int count1 = 3;
- int count2 = 2;
- // data1 indices: {0, 4, 8}
- // data2 indices: {0, 5}
-#pragma omp target update to(data1[0 : count1 : 4], data2[0 : count2 : 5])
-
- // Device modifies arrays
-#pragma omp target
- {
- for (int i = 0; i < len; i++) {
- data[i] += i; // becomes 2*i on device
- data1[i] += i; // becomes 2*i on device
- data2[i] += 100; // becomes i*10 + 100 on device
- }
- }
- }
-
- // Print results
- printf("from target array results (data):\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data[i]);
- printf("\n");
-
- printf("from target array results (data1, data2):\n");
- printf("data1:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data1[i]);
- printf("data2:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data2[i]);
- printf("\n");
-
- return 0;
-}
-
-// CHECK: from target array results (data):
-// CHECK: 0.000000
-// CHECK: 1.000000
-// CHECK: 2.000000
-// CHECK: 3.000000
-// CHECK: 4.000000
-// CHECK: 5.000000
-// CHECK: 6.000000
-// CHECK: 7.000000
-// CHECK: 8.000000
-// CHECK: 9.000000
-// CHECK: 10.000000
-// CHECK: 11.000000
-// CHECK: 12.000000
-// CHECK: 13.000000
-// CHECK: 14.000000
-// CHECK: 15.000000
-
-// CHECK: from target array results (data1, data2):
-// CHECK: data1:
-// CHECK: 0.000000
-// CHECK: 1.000000
-// CHECK: 2.000000
-// CHECK: 3.000000
-// CHECK: 4.000000
-// CHECK: 5.000000
-// CHECK: 6.000000
-// CHECK: 7.000000
-// CHECK: 8.000000
-// CHECK: 9.000000
-// CHECK: 10.000000
-// CHECK: 11.000000
-// CHECK: 12.000000
-// CHECK: 13.000000
-// CHECK: 14.000000
-// CHECK: 15.000000
-// CHECK: data2:
-// CHECK: 0.000000
-// CHECK: 10.000000
-// CHECK: 20.000000
-// CHECK: 30.000000
-// CHECK: 40.000000
-// CHECK: 50.000000
-// CHECK: 60.000000
-// CHECK: 70.000000
-// CHECK: 80.000000
-// CHECK: 90.000000
-// CHECK: 100.000000
-// CHECK: 110.000000
-// CHECK: 120.000000
-// CHECK: 130.000000
-// CHECK: 140.000000
-// CHECK: 150.000000
diff --git a/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c b/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c
deleted file mode 100644
index 1567e79814b84..0000000000000
--- a/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c
+++ /dev/null
@@ -1,43 +0,0 @@
-
-// RUN: %libomptarget-compile-run-and-check-generic
-#include <omp.h>
-#include <stdio.h>
-
-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;
- }
- }
-
-#pragma omp target data map(to : len, data[0 : len])
- {
-#pragma omp target
- for (int i = 0; i < len; i++) {
- data[i] += i;
- }
-
- int zero_count = 0;
-#pragma omp target update from(data[0 : zero_count : 2])
- }
-
- printf("from target array results:\n");
- for (int i = 0; i < len; i++)
- printf("%f\n", data[i]);
- printf("\n");
-
- return 0;
-}
-
-// CHECK: from target array results:
-// CHECK: 0.000000
-// CHECK: 1.000000
-// CHECK: 2.000000
-// CHECK: 3.000000
-// CHECK: 4.000000
-// CHECK: 5.000000
-// CHECK: 6.000000
-// CHECK: 7.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_count_expression_from.c b/offload/test/offloading/target_update_ptr_count_expression_from.c
deleted file mode 100644
index 2c5740b906f6c..0000000000000
--- a/offload/test/offloading/target_update_ptr_count_expression_from.c
+++ /dev/null
@@ -1,74 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-// This test checks that "update from" clause in OpenMP supports strided
-// sections. #pragma omp target update from(result[0:len/2:2]) updates every other
-// element from device
-#include <omp.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-#define N 32
-
-int main() {
- double *result = (double *)calloc(N, sizeof(double));
- int len = N;
-
- printf("initial host array values:\n");
- for (int i = 0; i < N; i++)
- printf("%f\n", result[i]);
- printf("\n");
-
-#pragma omp target data map(to : result[0 : N])
- {
-#pragma omp target map(alloc : result[0 : N])
- for (int i = 0; i < N; i++)
- result[i] += i;
-
- // Update strided elements from device: even indices 0,2,4,...,30
-#pragma omp target update from(result[0 : len/2 : 2])
- }
-
- printf("after target update from (even indices up to 30 updated):\n");
- for (int i = 0; i < N; i++)
- printf("%f\n", result[i]);
- printf("\n");
-
- // Expected: even indices i, odd indices 0
- // CHECK: 0.000000
- // CHECK: 0.000000
- // CHECK: 2.000000
- // CHECK: 0.000000
- // CHECK: 4.000000
- // CHECK: 0.000000
- // CHECK: 6.000000
- // CHECK: 0.000000
- // CHECK: 8.000000
- // CHECK: 0.000000
- // CHECK: 10.000000
- // CHECK: 0.000000
- // CHECK: 12.000000
- // CHECK: 0.000000
- // CHECK: 14.000000
- // CHECK: 0.000000
- // CHECK: 16.000000
- // CHECK: 0.000000
- // CHECK: 18.000000
- // CHECK: 0.000000
- // CHECK: 20.000000
- // CHECK: 0.000000
- // CHECK: 22.000000
- // CHECK: 0.000000
- // CHECK: 24.000000
- // CHECK: 0.000000
- // CHECK: 26.000000
- // CHECK: 0.000000
- // CHECK: 28.000000
- // CHECK: 0.000000
- // CHECK: 30.000000
- // CHECK: 0.000000
- // CHECK-NOT: 1.000000
- // CHECK-NOT: 3.000000
- // CHECK-NOT: 31.000000
-
- free(result);
- return 0;
-}
diff --git a/offload/test/offloading/target_update_ptr_count_expression_to.c b/offload/test/offloading/target_update_ptr_count_expression_to.c
deleted file mode 100644
index 3900592f10f93..0000000000000
--- a/offload/test/offloading/target_update_ptr_count_expression_to.c
+++ /dev/null
@@ -1,82 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-// This test checks that "update to" clause in OpenMP supports strided sections.
-// #pragma omp target update to(result[0:len/2:2]) updates every other element
-// (stride 2)
-
-#include <omp.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-#define N 16
-
-int main() {
- double *result = (double *)calloc(N, sizeof(double));
- int len = N;
-
- // Initialize on host
- for (int i = 0; i < N; i++) {
- result[i] = i;
- }
-
- // Initial values
- printf("original host array values:\n");
- for (int i = 0; i < N; i++)
- printf("%f\n", result[i]);
- printf("\n");
-
-#pragma omp target data map(tofrom : result[0 : N])
- {
-// Update strided elements to device: indices 0,2,4,6
-#pragma omp target update to(result[0 : len/2 : 2])
-
-#pragma omp target
- {
- for (int i = 0; i < N; i++) {
- result[i] += i;
- }
- }
- }
-
- printf("from target array results:\n");
- for (int i = 0; i < N; i++)
- printf("%f\n", result[i]);
-
- // 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-NEXT: 10.000000
- // CHECK-NEXT: 11.000000
- // CHECK-NEXT: 12.000000
- // CHECK-NEXT: 13.000000
- // CHECK-NEXT: 14.000000
- // CHECK-NEXT: 15.000000
-
- // CHECK: from target array results:
- // CHECK-NEXT: 0.000000
- // CHECK-NEXT: 2.000000
- // CHECK-NEXT: 4.000000
- // CHECK-NEXT: 6.000000
- // CHECK-NEXT: 8.000000
- // CHECK-NEXT: 10.000000
- // CHECK-NEXT: 12.000000
- // CHECK-NEXT: 14.000000
- // CHECK-NEXT: 16.000000
- // CHECK-NEXT: 18.000000
- // CHECK-NEXT: 20.000000
- // CHECK-NEXT: 22.000000
- // CHECK-NEXT: 24.000000
- // CHECK-NEXT: 26.000000
- // CHECK-NEXT: 28.000000
- // CHECK-NEXT: 30.000000
-
- free(result);
- return 0;
-}
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_count_expression_from.c b/offload/test/offloading/target_update_strided_struct_count_expression_from.c
deleted file mode 100644
index bb3e2ba577f7f..0000000000000
--- a/offload/test/offloading/target_update_strided_struct_count_expression_from.c
+++ /dev/null
@@ -1,86 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-// This test checks that "update from" with user-defined mapper supports strided
-// sections using fixed-size arrays in structs.
-
-#include <omp.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-#define N 16
-
-typedef struct {
- double data[N];
- size_t 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;
-
- 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("%f\n", s.data[i]);
- printf("\n");
-
-#pragma omp target data map(mapper(custom), tofrom : s)
- {
-// Execute on device with explicit mapper
-#pragma omp target map(mapper(custom), tofrom : s)
- {
- for (int i = 0; i < s.len; i++) {
- s.data[i] += i;
- }
- }
-
-// Update strided elements from device: indices 0,2,4,6,8,10,12,14
-#pragma omp target update from(s.data[0 : s.len/2 : 2])
- }
-
- printf("from target array results:\n");
- for (int i = 0; i < N; i++)
- printf("%f\n", s.data[i]);
-
- // 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-NEXT: 10.000000
- // CHECK-NEXT: 11.000000
- // CHECK-NEXT: 12.000000
- // CHECK-NEXT: 13.000000
- // CHECK-NEXT: 14.000000
- // CHECK-NEXT: 15.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-NEXT: 20.000000
- // CHECK-NEXT: 11.000000
- // CHECK-NEXT: 24.000000
- // CHECK-NEXT: 13.000000
- // CHECK-NEXT: 28.000000
- // CHECK-NEXT: 15.000000
-
- return 0;
-}
diff --git a/offload/test/offloading/target_update_strided_struct_count_expression_to.c b/offload/test/offloading/target_update_strided_struct_count_expression_to.c
deleted file mode 100644
index c4bc7faf367a0..0000000000000
--- a/offload/test/offloading/target_update_strided_struct_count_expression_to.c
+++ /dev/null
@@ -1,98 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-// This test checks that "update to" with struct member arrays supports strided
-// sections using fixed-size arrays in structs.
-
-#include <omp.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-#define N 16
-
-typedef struct {
- double data[N];
- int len;
-} T;
-
-int main() {
- T s;
- s.len = N;
-
- // Initialize struct array on host with simple sequential values
- for (int i = 0; i < N; i++) {
- s.data[i] = i;
- }
-
- printf("original host struct 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)
- {
-// Initialize device struct array to 20
-#pragma omp target map(tofrom : s)
- {
- for (int i = 0; i < s.len; i++) {
- s.data[i] = 20.0;
- }
- }
-
- // Modify host struct data for strided elements (set to 10)
- for (int i = 0; i < 8; i++) {
- s.data[i * 2] = 10.0; // Set even indices to 10
- }
-
-// indices 0,2,4,6,8,10,12,14
-#pragma omp target update to(s.data[0 : s.len/2 : 2])
-
-// Execute on device - add 5 to verify update worked
-#pragma omp target map(tofrom : s)
- {
- for (int i = 0; i < s.len; i++) {
- s.data[i] += 5.0;
- }
- }
- }
-
- printf("after target update to struct:\n");
- for (int i = 0; i < N; i++)
- printf("%.1f\n", s.data[i]);
-
- // CHECK: original host struct 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: after target update to struct:
- // 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
-
- return 0;
-}
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 4b4d6c4eb5c1b12a9c60d1f755bfa82c8f21f9d5 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 19 Jan 2026 02:52:03 -0500
Subject: [PATCH 5/5] Apply minimal formatting
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 297 +++++++++++++++++++-------
1 file changed, 220 insertions(+), 77 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index bd7189dbf23a4..e6f5f00a86922 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -492,11 +492,11 @@ enum OpenMPLocationFlags : unsigned {
/// member */
/// kmp_int32 reserved_2; /**< not really used in Fortran any more;
/// see above */
-/// #if USE_ITT_BUILD
+///#if USE_ITT_BUILD
/// /* but currently used for storing
/// region-specific ITT */
/// /* contextual information. */
-/// #endif /* USE_ITT_BUILD */
+///#endif /* USE_ITT_BUILD */
/// kmp_int32 reserved_3; /**< source[4] in Fortran, do not use for
/// C++ */
/// char const *psource; /**< String describing the source location.
@@ -716,16 +716,16 @@ static void EmitOMPAggregateInit(CodeGenFunction &CGF, Address DestAddr,
if (DRD) {
// Shift the address forward by one element.
- llvm::Value *SrcElementNext =
- CGF.Builder.CreateConstGEP1_32(SrcAddr.getElementType(), SrcElementPHI,
- /*Idx0=*/1, "omp.arraycpy.dest.element");
+ llvm::Value *SrcElementNext = CGF.Builder.CreateConstGEP1_32(
+ SrcAddr.getElementType(), SrcElementPHI, /*Idx0=*/1,
+ "omp.arraycpy.dest.element");
SrcElementPHI->addIncoming(SrcElementNext, CGF.Builder.GetInsertBlock());
}
// Shift the address forward by one element.
- llvm::Value *DestElementNext =
- CGF.Builder.CreateConstGEP1_32(DestAddr.getElementType(), DestElementPHI,
- /*Idx0=*/1, "omp.arraycpy.dest.element");
+ llvm::Value *DestElementNext = CGF.Builder.CreateConstGEP1_32(
+ DestAddr.getElementType(), DestElementPHI, /*Idx0=*/1,
+ "omp.arraycpy.dest.element");
// Check whether we've reached the end.
llvm::Value *Done =
CGF.Builder.CreateICmpEQ(DestElementNext, DestEnd, "omp.arraycpy.done");
@@ -975,8 +975,8 @@ Address ReductionCodeGen::adjustPrivateAddress(CodeGenFunction &CGF, unsigned N,
llvm::Value *PrivatePointer =
CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
PrivateAddr.emitRawPointer(CGF), SharedAddr.getType());
- llvm::Value *Ptr = CGF.Builder.CreateGEP(SharedAddr.getElementType(),
- PrivatePointer, Adjustment);
+ llvm::Value *Ptr = CGF.Builder.CreateGEP(
+ SharedAddr.getElementType(), PrivatePointer, Adjustment);
return castToBase(CGF, OrigVD->getType(),
SharedAddresses[N].first.getType(),
OriginalBaseLValue.getAddress(), Ptr);
@@ -1605,11 +1605,12 @@ Address CGOpenMPRuntime::getAddrOfThreadPrivate(CodeGenFunction &CGF,
CGF.Builder.CreatePointerCast(VDAddr.emitRawPointer(CGF), CGM.Int8PtrTy),
CGM.getSize(CGM.GetTargetTypeStoreSize(VarTy)),
getOrCreateThreadPrivateCache(VD)};
- return Address(CGF.EmitRuntimeCall(
- OMPBuilder.getOrCreateRuntimeFunction(
- CGM.getModule(), OMPRTL___kmpc_threadprivate_cached),
- Args),
- CGF.Int8Ty, VDAddr.getAlignment());
+ return Address(
+ CGF.EmitRuntimeCall(
+ OMPBuilder.getOrCreateRuntimeFunction(
+ CGM.getModule(), OMPRTL___kmpc_threadprivate_cached),
+ Args),
+ CGF.Int8Ty, VDAddr.getAlignment());
}
void CGOpenMPRuntime::emitThreadPrivateVarInit(
@@ -1634,8 +1635,8 @@ void CGOpenMPRuntime::emitThreadPrivateVarInit(
}
llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition(
- const VarDecl *VD, Address VDAddr, SourceLocation Loc, bool PerformInit,
- CodeGenFunction *CGF) {
+ const VarDecl *VD, Address VDAddr, SourceLocation Loc,
+ bool PerformInit, CodeGenFunction *CGF) {
if (CGM.getLangOpts().OpenMPUseTLS &&
CGM.getContext().getTargetInfo().isTLSSupported())
return nullptr;
@@ -1697,8 +1698,7 @@ llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition(
auto NL = ApplyDebugLocation::CreateEmpty(DtorCGF);
DtorCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, Fn, FI, Args,
Loc, Loc);
- // Create a scope with an artificial location for the body of this
- // function.
+ // Create a scope with an artificial location for the body of this function.
auto AL = ApplyDebugLocation::CreateArtificial(DtorCGF);
llvm::Value *ArgVal = DtorCGF.EmitLoadOfScalar(
DtorCGF.GetAddrOfLocalVar(&Dst),
@@ -1776,12 +1776,126 @@ void CGOpenMPRuntime::emitDeclareTargetFunction(const FunctionDecl *FD,
Addr->setVisibility(llvm::GlobalValue::ProtectedVisibility);
}
+ // Register the indirect Vtable:
+ // This is similar to OMPTargetGlobalVarEntryIndirect, except that the
+ // size field refers to the size of memory pointed to, not the size of
+ // the pointer symbol itself (which is implicitly the size of a pointer).
OMPBuilder.OffloadInfoManager.registerDeviceGlobalVarEntryInfo(
Name, Addr, CGM.GetTargetTypeStoreSize(CGM.VoidPtrTy).getQuantity(),
llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect,
llvm::GlobalValue::WeakODRLinkage);
}
+void CGOpenMPRuntime::registerVTableOffloadEntry(llvm::GlobalVariable *VTable,
+ const VarDecl *VD) {
+ // TODO: add logic to avoid duplicate vtable registrations per
+ // translation unit; though for external linkage, this should no
+ // longer be an issue - or at least we can avoid the issue by
+ // checking for an existing offloading entry. But, perhaps the
+ // better approach is to defer emission of the vtables and offload
+ // entries until later (by tracking a list of items that need to be
+ // emitted).
+
+ llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
+
+ // Generate a new externally visible global to point to the
+ // internally visible vtable. Doing this allows us to keep the
+ // visibility and linkage of the associated vtable unchanged while
+ // allowing the runtime to access its value. The externally
+ // visible global var needs to be emitted with a unique mangled
+ // name that won't conflict with similarly named (internal)
+ // vtables in other translation units.
+
+ // Register vtable with source location of dynamic object in map
+ // clause.
+ llvm::TargetRegionEntryInfo EntryInfo = getEntryInfoFromPresumedLoc(
+ CGM, OMPBuilder, VD->getCanonicalDecl()->getBeginLoc(),
+ VTable->getName());
+
+ llvm::GlobalVariable *Addr = VTable;
+ SmallString<128> AddrName;
+ OMPBuilder.OffloadInfoManager.getTargetRegionEntryFnName(AddrName, EntryInfo);
+ AddrName.append("addr");
+
+ if (CGM.getLangOpts().OpenMPIsTargetDevice) {
+ Addr = new llvm::GlobalVariable(
+ CGM.getModule(), VTable->getType(),
+ /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, VTable,
+ AddrName,
+ /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
+ CGM.getModule().getDataLayout().getDefaultGlobalsAddressSpace());
+ Addr->setVisibility(llvm::GlobalValue::ProtectedVisibility);
+ }
+ OMPBuilder.OffloadInfoManager.registerDeviceGlobalVarEntryInfo(
+ AddrName, VTable,
+ CGM.getDataLayout().getTypeAllocSize(VTable->getInitializer()->getType()),
+ llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable,
+ llvm::GlobalValue::WeakODRLinkage);
+}
+
+void CGOpenMPRuntime::emitAndRegisterVTable(CodeGenModule &CGM,
+ CXXRecordDecl *CXXRecord,
+ const VarDecl *VD) {
+ // Register C++ VTable to OpenMP Offload Entry if it's a new
+ // CXXRecordDecl.
+ if (CXXRecord && CXXRecord->isDynamicClass() &&
+ !CGM.getOpenMPRuntime().VTableDeclMap.contains(CXXRecord)) {
+ auto Res = CGM.getOpenMPRuntime().VTableDeclMap.try_emplace(CXXRecord, VD);
+ if (Res.second) {
+ CGM.EmitVTable(CXXRecord);
+ CodeGenVTables VTables = CGM.getVTables();
+ llvm::GlobalVariable *VTablesAddr = VTables.GetAddrOfVTable(CXXRecord);
+ assert(VTablesAddr && "Expected non-null VTable address");
+ CGM.getOpenMPRuntime().registerVTableOffloadEntry(VTablesAddr, VD);
+ // Emit VTable for all the fields containing dynamic CXXRecord
+ for (const FieldDecl *Field : CXXRecord->fields()) {
+ if (CXXRecordDecl *RecordDecl = Field->getType()->getAsCXXRecordDecl())
+ emitAndRegisterVTable(CGM, RecordDecl, VD);
+ }
+ // Emit VTable for all dynamic parent class
+ for (CXXBaseSpecifier &Base : CXXRecord->bases()) {
+ if (CXXRecordDecl *BaseDecl = Base.getType()->getAsCXXRecordDecl())
+ emitAndRegisterVTable(CGM, BaseDecl, VD);
+ }
+ }
+ }
+}
+
+void CGOpenMPRuntime::registerVTable(const OMPExecutableDirective &D) {
+ // Register VTable by scanning through the map clause of OpenMP target region.
+ // Get CXXRecordDecl and VarDecl from Expr.
+ auto GetVTableDecl = [](const Expr *E) {
+ QualType VDTy = E->getType();
+ CXXRecordDecl *CXXRecord = nullptr;
+ if (const auto *RefType = VDTy->getAs<LValueReferenceType>())
+ VDTy = RefType->getPointeeType();
+ if (VDTy->isPointerType())
+ CXXRecord = VDTy->getPointeeType()->getAsCXXRecordDecl();
+ else
+ CXXRecord = VDTy->getAsCXXRecordDecl();
+
+ const VarDecl *VD = nullptr;
+ if (auto *DRE = dyn_cast<DeclRefExpr>(E)) {
+ VD = cast<VarDecl>(DRE->getDecl());
+ } else if (auto *MRE = dyn_cast<MemberExpr>(E)) {
+ if (auto *BaseDRE = dyn_cast<DeclRefExpr>(MRE->getBase())) {
+ if (auto *BaseVD = dyn_cast<VarDecl>(BaseDRE->getDecl()))
+ VD = BaseVD;
+ }
+ }
+ return std::pair<CXXRecordDecl *, const VarDecl *>(CXXRecord, VD);
+ };
+ // Collect VTable from OpenMP map clause.
+ for (const auto *C : D.getClausesOfKind<OMPMapClause>()) {
+ for (const auto *E : C->varlist()) {
+ auto DeclPair = GetVTableDecl(E);
+ // Ensure VD is not null
+ if (DeclPair.second)
+ emitAndRegisterVTable(CGM, DeclPair.first, DeclPair.second);
+ }
+ }
+}
+
Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF,
QualType VarType,
StringRef Name) {
@@ -1942,7 +2056,8 @@ Address CGOpenMPRuntime::emitThreadIDAddress(CodeGenFunction &CGF,
QualType Int32Ty =
CGF.getContext().getIntTypeForBitwidth(/*DestWidth*/ 32, /*Signed*/ true);
Address ThreadIDTemp = CGF.CreateMemTemp(Int32Ty, /*Name*/ ".threadid_temp.");
- CGF.EmitStoreOfScalar(ThreadID, CGF.MakeAddrLValue(ThreadIDTemp, Int32Ty));
+ CGF.EmitStoreOfScalar(ThreadID,
+ CGF.MakeAddrLValue(ThreadIDTemp, Int32Ty));
return ThreadIDTemp;
}
@@ -2449,8 +2564,8 @@ bool CGOpenMPRuntime::isStaticChunked(OpenMPScheduleClauseKind ScheduleKind,
return Schedule == OMP_sch_static_chunked;
}
-bool CGOpenMPRuntime::isStaticChunked(OpenMPDistScheduleClauseKind ScheduleKind,
- bool Chunked) const {
+bool CGOpenMPRuntime::isStaticChunked(
+ OpenMPDistScheduleClauseKind ScheduleKind, bool Chunked) const {
OpenMPSchedType Schedule = getRuntimeSchedule(ScheduleKind, Chunked);
return Schedule == OMP_dist_sch_static_chunked;
}
@@ -2614,10 +2729,10 @@ void CGOpenMPRuntime::emitForStaticInit(CodeGenFunction &CGF,
ScheduleKind.Schedule, Values.Chunk != nullptr, Values.Ordered);
assert((isOpenMPWorksharingDirective(DKind) || (DKind == OMPD_loop)) &&
"Expected loop-based or sections-based directive.");
- llvm::Value *UpdatedLocation = emitUpdateLocation(
- CGF, Loc,
- isOpenMPLoopDirective(DKind) ? OMP_IDENT_WORK_LOOP
- : OMP_IDENT_WORK_SECTIONS);
+ llvm::Value *UpdatedLocation = emitUpdateLocation(CGF, Loc,
+ isOpenMPLoopDirective(DKind)
+ ? OMP_IDENT_WORK_LOOP
+ : OMP_IDENT_WORK_SECTIONS);
llvm::Value *ThreadId = getThreadID(CGF, Loc);
llvm::FunctionCallee StaticInitFunction =
OMPBuilder.createForStaticInitFunction(Values.IVSize, Values.IVSigned,
@@ -2692,8 +2807,9 @@ void CGOpenMPRuntime::emitForOrderedIterationEnd(CodeGenFunction &CGF,
llvm::Value *CGOpenMPRuntime::emitForNext(CodeGenFunction &CGF,
SourceLocation Loc, unsigned IVSize,
- bool IVSigned, Address IL, Address LB,
- Address UB, Address ST) {
+ bool IVSigned, Address IL,
+ Address LB, Address UB,
+ Address ST) {
// Call __kmpc_dispatch_next(
// ident_t *loc, kmp_int32 tid, kmp_int32 *p_lastiter,
// kmp_int[32|64] *p_lower, kmp_int[32|64] *p_upper,
@@ -2843,6 +2959,13 @@ void CGOpenMPRuntime::createOffloadEntriesAndInfoMetadata() {
case llvm::OpenMPIRBuilder::EMIT_MD_GLOBAL_VAR_LINK_ERROR: {
CGM.getDiags().Report(diag::err_target_var_offloading_entry_incorrect);
} break;
+ case llvm::OpenMPIRBuilder::EMIT_MD_GLOBAL_VAR_INDIRECT_ERROR: {
+ unsigned DiagID = CGM.getDiags().getCustomDiagID(
+ DiagnosticsEngine::Error, "Offloading entry for indirect declare "
+ "target variable is incorrect: the "
+ "address is invalid.");
+ CGM.getDiags().Report(DiagID);
+ } break;
}
};
@@ -2889,8 +3012,8 @@ static bool isAllocatableDecl(const VarDecl *VD) {
!AA->getAllocator());
}
-static RecordDecl *createPrivatesRecordDecl(CodeGenModule &CGM,
- ArrayRef<PrivateDataTy> Privates) {
+static RecordDecl *
+createPrivatesRecordDecl(CodeGenModule &CGM, ArrayRef<PrivateDataTy> Privates) {
if (!Privates.empty()) {
ASTContext &C = CGM.getContext();
// Build struct .kmp_privates_t. {
@@ -3394,6 +3517,7 @@ static bool checkInitIsRequired(CodeGenFunction &CGF,
return InitRequired;
}
+
/// Emit task_dup function (for initialization of
/// private/firstprivate/lastprivate vars and last_iter flag)
/// \code
@@ -3731,6 +3855,7 @@ CGOpenMPRuntime::emitTaskInit(CodeGenFunction &CGF, SourceLocation Loc,
PriorityFlag = 0x20,
DetachableFlag = 0x40,
FreeAgentFlag = 0x80,
+ TransparentFlag = 0x100,
};
unsigned Flags = Data.Tied ? TiedFlag : 0;
bool NeedsCleanup = false;
@@ -3745,6 +3870,9 @@ CGOpenMPRuntime::emitTaskInit(CodeGenFunction &CGF, SourceLocation Loc,
if (Kind == OMPC_THREADSET_omp_pool)
Flags = Flags | FreeAgentFlag;
}
+ if (D.getSingleClause<OMPTransparentClause>())
+ Flags |= TransparentFlag;
+
if (Data.Priority.getInt())
Flags = Flags | PriorityFlag;
if (D.hasClausesOfKind<OMPDetachClause>())
@@ -3757,14 +3885,10 @@ CGOpenMPRuntime::emitTaskInit(CodeGenFunction &CGF, SourceLocation Loc,
: CGF.Builder.getInt32(Data.Final.getInt() ? FinalFlag : 0);
TaskFlags = CGF.Builder.CreateOr(TaskFlags, CGF.Builder.getInt32(Flags));
llvm::Value *SharedsSize = CGM.getSize(C.getTypeSizeInChars(SharedsTy));
- SmallVector<llvm::Value *, 8> AllocArgs = {
- emitUpdateLocation(CGF, Loc),
- getThreadID(CGF, Loc),
- TaskFlags,
- KmpTaskTWithPrivatesTySize,
- SharedsSize,
- CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(TaskEntry,
- KmpRoutineEntryPtrTy)};
+ SmallVector<llvm::Value *, 8> AllocArgs = {emitUpdateLocation(CGF, Loc),
+ getThreadID(CGF, Loc), TaskFlags, KmpTaskTWithPrivatesTySize,
+ SharedsSize, CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+ TaskEntry, KmpRoutineEntryPtrTy)};
llvm::Value *NewTask;
if (D.hasClausesOfKind<OMPNowaitClause>()) {
// Check if we have any device clause associated with the directive.
@@ -3955,13 +4079,13 @@ CGOpenMPRuntime::emitTaskInit(CodeGenFunction &CGF, SourceLocation Loc,
// Copy shareds if there are any.
Address KmpTaskSharedsPtr = Address::invalid();
if (!SharedsTy->castAsRecordDecl()->field_empty()) {
- KmpTaskSharedsPtr =
- Address(CGF.EmitLoadOfScalar(
- CGF.EmitLValueForField(
- TDBase, *std::next(KmpTaskTQTyRD->field_begin(),
- KmpTaskTShareds)),
- Loc),
- CGF.Int8Ty, CGM.getNaturalTypeAlignment(SharedsTy));
+ KmpTaskSharedsPtr = Address(
+ CGF.EmitLoadOfScalar(
+ CGF.EmitLValueForField(
+ TDBase,
+ *std::next(KmpTaskTQTyRD->field_begin(), KmpTaskTShareds)),
+ Loc),
+ CGF.Int8Ty, CGM.getNaturalTypeAlignment(SharedsTy));
LValue Dest = CGF.MakeAddrLValue(KmpTaskSharedsPtr, SharedsTy);
LValue Src = CGF.MakeAddrLValue(Shareds, SharedsTy);
CGF.EmitAggregateCopy(Dest, Src, SharedsTy, AggValueSlot::DoesNotOverlap);
@@ -4544,7 +4668,7 @@ void CGOpenMPRuntime::emitTaskCall(CodeGenFunction &CGF, SourceLocation Loc,
// list is not empty
llvm::Value *ThreadID = getThreadID(CGF, Loc);
llvm::Value *UpLoc = emitUpdateLocation(CGF, Loc);
- llvm::Value *TaskArgs[] = {UpLoc, ThreadID, NewTask};
+ llvm::Value *TaskArgs[] = { UpLoc, ThreadID, NewTask };
llvm::Value *DepTaskArgs[7];
if (!Data.Dependences.empty()) {
DepTaskArgs[0] = UpLoc;
@@ -4789,12 +4913,12 @@ static void EmitOMPAggregateReduction(
Scope.ForceCleanup();
// Shift the address forward by one element.
- llvm::Value *LHSElementNext =
- CGF.Builder.CreateConstGEP1_32(LHSAddr.getElementType(), LHSElementPHI,
- /*Idx0=*/1, "omp.arraycpy.dest.element");
- llvm::Value *RHSElementNext =
- CGF.Builder.CreateConstGEP1_32(RHSAddr.getElementType(), RHSElementPHI,
- /*Idx0=*/1, "omp.arraycpy.src.element");
+ llvm::Value *LHSElementNext = CGF.Builder.CreateConstGEP1_32(
+ LHSAddr.getElementType(), LHSElementPHI, /*Idx0=*/1,
+ "omp.arraycpy.dest.element");
+ llvm::Value *RHSElementNext = CGF.Builder.CreateConstGEP1_32(
+ RHSAddr.getElementType(), RHSElementPHI, /*Idx0=*/1,
+ "omp.arraycpy.src.element");
// Check whether we've reached the end.
llvm::Value *Done =
CGF.Builder.CreateICmpEQ(LHSElementNext, LHSEnd, "omp.arraycpy.done");
@@ -5744,7 +5868,7 @@ llvm::Value *CGOpenMPRuntime::emitTaskReductionInit(
const FieldDecl *SharedFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *OrigFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *SizeFD = addFieldToRecordDecl(C, RD, C.getSizeType());
- const FieldDecl *InitFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
+ const FieldDecl *InitFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *FiniFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *CombFD = addFieldToRecordDecl(C, RD, C.VoidPtrTy);
const FieldDecl *FlagsFD = addFieldToRecordDecl(
@@ -6250,11 +6374,12 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
CGM.handleAMDGPUWavesPerEUAttr(OutlinedFn, Attr);
}
}
+ registerVTable(D);
}
/// Checks if the expression is constant or does not have non-trivial function
/// calls.
-static bool isTrivial(ASTContext &Ctx, const Expr *E) {
+static bool isTrivial(ASTContext &Ctx, const Expr * E) {
// We can skip constant expressions.
// We can skip expressions with trivial calls or simple expressions.
return (E->isEvaluatable(Ctx, Expr::SE_AllowUndefinedBehavior) ||
@@ -6449,11 +6574,10 @@ llvm::Value *CGOpenMPRuntime::emitNumTeamsForTargetDirective(
const auto *CS = D.getInnermostCapturedStmt();
CGOpenMPInnerExprInfo CGInfo(CGF, *CS);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
- llvm::Value *NumTeamsVal =
- CGF.EmitScalarExpr(NumTeams,
- /*IgnoreResultAssign*/ true);
+ llvm::Value *NumTeamsVal = CGF.EmitScalarExpr(NumTeams,
+ /*IgnoreResultAssign*/ true);
return Bld.CreateIntCast(NumTeamsVal, CGF.Int32Ty,
- /*isSigned=*/true);
+ /*isSigned=*/true);
}
case OMPD_target_teams:
case OMPD_target_teams_distribute:
@@ -6461,11 +6585,10 @@ llvm::Value *CGOpenMPRuntime::emitNumTeamsForTargetDirective(
case OMPD_target_teams_distribute_parallel_for:
case OMPD_target_teams_distribute_parallel_for_simd: {
CodeGenFunction::RunCleanupsScope NumTeamsScope(CGF);
- llvm::Value *NumTeamsVal =
- CGF.EmitScalarExpr(NumTeams,
- /*IgnoreResultAssign*/ true);
+ llvm::Value *NumTeamsVal = CGF.EmitScalarExpr(NumTeams,
+ /*IgnoreResultAssign*/ true);
return Bld.CreateIntCast(NumTeamsVal, CGF.Int32Ty,
- /*isSigned=*/true);
+ /*isSigned=*/true);
}
default:
break;
@@ -8132,8 +8255,8 @@ class MappableExprsHandler {
LowestElem, CGF.VoidPtrTy, CGF.Int8Ty),
TypeSize.getQuantity() - 1);
PartialStruct.HighestElem = {
- std::numeric_limits<
- decltype(PartialStruct.HighestElem.first)>::max(),
+ std::numeric_limits<decltype(
+ PartialStruct.HighestElem.first)>::max(),
HB};
PartialStruct.Base = BP;
PartialStruct.LB = LB;
@@ -10922,6 +11045,17 @@ void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
if (!S)
return;
+ // Register vtable from device for target data and target directives.
+ // Add this block here since scanForTargetRegionsFunctions ignores
+ // target data by checking if S is a executable directive (target).
+ if (auto *E = dyn_cast<OMPExecutableDirective>(S);
+ E && isOpenMPTargetDataManagementDirective(E->getDirectiveKind())) {
+ // Don't need to check if it's device compile
+ // since scanForTargetRegionsFunctions currently only called
+ // in device compilation.
+ registerVTable(*E);
+ }
+
// Codegen OMP target directives that offload compute to the device.
bool RequiresDeviceCodegen =
isa<OMPExecutableDirective>(S) &&
@@ -11267,7 +11401,7 @@ bool CGOpenMPRuntime::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
return false;
const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
- switch (A->getAllocatorType()) {
+ switch(A->getAllocatorType()) {
case OMPAllocateDeclAttr::OMPNullMemAlloc:
case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
// Not supported, fallback to the default mem space.
@@ -11371,8 +11505,7 @@ void CGOpenMPRuntime::emitNumTeamsClause(CodeGenFunction &CGF,
CGF.CGM.Int32Ty, /* isSigned = */ true)
: CGF.Builder.getInt32(0);
- // Build call __kmpc_push_num_teamss(&loc, global_tid, num_teams,
- // thread_limit)
+ // Build call __kmpc_push_num_teamss(&loc, global_tid, num_teams, thread_limit)
llvm::Value *PushNumTeamsArgs[] = {RTLoc, getThreadID(CGF, Loc), NumTeamsVal,
ThreadLimitVal};
CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
@@ -11681,7 +11814,7 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
}
namespace {
-/// Kind of parameter in a function with 'declare simd' directive.
+ /// Kind of parameter in a function with 'declare simd' directive.
enum ParamKindTy {
Linear,
LinearRef,
@@ -11807,10 +11940,18 @@ emitX86DeclareSimdFunction(const FunctionDecl *FD, llvm::Function *Fn,
unsigned VecRegSize;
};
ISADataTy ISAData[] = {
- {'b', 128}, // SSE
- {'c', 256}, // AVX
- {'d', 256}, // AVX2
- {'e', 512}, // AVX512
+ {
+ 'b', 128
+ }, // SSE
+ {
+ 'c', 256
+ }, // AVX
+ {
+ 'd', 256
+ }, // AVX2
+ {
+ 'e', 512
+ }, // AVX512
};
llvm::SmallVector<char, 2> Masked;
switch (State) {
@@ -12791,8 +12932,7 @@ Address CGOpenMPRuntime::emitLastprivateConditionalInit(CodeGenFunction &CGF,
FiredField = addFieldToRecordDecl(C, RD, C.CharTy);
RD->completeDefinition();
NewType = C.getCanonicalTagType(RD);
- Address Addr =
- CGF.CreateMemTemp(NewType, C.getDeclAlign(VD), VD->getName());
+ Address Addr = CGF.CreateMemTemp(NewType, C.getDeclAlign(VD), VD->getName());
BaseLVal = CGF.MakeAddrLValue(Addr, NewType, AlignmentSource::Decl);
I->getSecond().try_emplace(VD, NewType, VDField, FiredField, BaseLVal);
} else {
@@ -12801,7 +12941,8 @@ Address CGOpenMPRuntime::emitLastprivateConditionalInit(CodeGenFunction &CGF,
FiredField = std::get<2>(VI->getSecond());
BaseLVal = std::get<3>(VI->getSecond());
}
- LValue FiredLVal = CGF.EmitLValueForField(BaseLVal, FiredField);
+ LValue FiredLVal =
+ CGF.EmitLValueForField(BaseLVal, FiredField);
CGF.EmitStoreOfScalar(
llvm::ConstantInt::getNullValue(CGF.ConvertTypeForMem(C.CharTy)),
FiredLVal);
@@ -12988,7 +13129,7 @@ void CGOpenMPRuntime::checkAndEmitLastprivateConditional(CodeGenFunction &CGF,
assert(It != LastprivateConditionalToTypes[FoundFn].end() &&
"Lastprivate conditional is not found in outer region.");
QualType StructTy = std::get<0>(It->getSecond());
- const FieldDecl *FiredDecl = std::get<2>(It->getSecond());
+ const FieldDecl* FiredDecl = std::get<2>(It->getSecond());
LValue PrivLVal = CGF.EmitLValue(FoundE);
Address StructAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
PrivLVal.getAddress(),
@@ -13342,7 +13483,9 @@ bool CGOpenMPSIMDRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
llvm_unreachable("Not supported in SIMD-only mode");
}
-bool CGOpenMPSIMDRuntime::emitTargetGlobal(GlobalDecl GD) { return false; }
+bool CGOpenMPSIMDRuntime::emitTargetGlobal(GlobalDecl GD) {
+ return false;
+}
void CGOpenMPSIMDRuntime::emitTeamsCall(CodeGenFunction &CGF,
const OMPExecutableDirective &D,
More information about the cfe-commits
mailing list