[clang] [llvm] [Offload][Tests] Non-contiguous_update_to_tests (PR #169623)
Amit Tiwari via llvm-commits
llvm-commits at lists.llvm.org
Mon Dec 1 00:20:27 PST 2025
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/169623
>From e38e0feaf7bb13e6147eca002cd616499253c595 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Wed, 26 Nov 2025 03:58:44 -0500
Subject: [PATCH] non-contiguous_update_to_tests
---
... => target_update_strided_messages_from.c} | 5 -
.../target_update_strided_messages_to.c | 33 +++++
...t_update_strided_multiple_messages_from.c} | 23 ----
...rget_update_strided_multiple_messages_to.c | 25 ++++
...et_update_strided_partial_messages_from.c} | 11 +-
...arget_update_strided_partial_messages_to.c | 25 ++++
...pdate.c => strided_multiple_update_from.c} | 0
.../offloading/strided_multiple_update_to.c | 122 ++++++++++++++++++
...update.c => strided_partial_update_from.c} | 0
.../offloading/strided_partial_update_to.c | 72 +++++++++++
...strided_update.c => strided_update_from.c} | 0
offload/test/offloading/strided_update_to.c | 72 +++++++++++
12 files changed, 351 insertions(+), 37 deletions(-)
rename clang/test/OpenMP/{target_update_strided_messages.c => target_update_strided_messages_from.c} (92%)
create mode 100644 clang/test/OpenMP/target_update_strided_messages_to.c
rename clang/test/OpenMP/{target_update_strided_multiple_messages.c => target_update_strided_multiple_messages_from.c} (60%)
create mode 100644 clang/test/OpenMP/target_update_strided_multiple_messages_to.c
rename clang/test/OpenMP/{target_update_strided_partial_messages.c => target_update_strided_partial_messages_from.c} (59%)
create mode 100644 clang/test/OpenMP/target_update_strided_partial_messages_to.c
rename offload/test/offloading/{strided_multiple_update.c => strided_multiple_update_from.c} (100%)
create mode 100644 offload/test/offloading/strided_multiple_update_to.c
rename offload/test/offloading/{strided_partial_update.c => strided_partial_update_from.c} (100%)
create mode 100644 offload/test/offloading/strided_partial_update_to.c
rename offload/test/offloading/{strided_update.c => strided_update_from.c} (100%)
create mode 100644 offload/test/offloading/strided_update_to.c
diff --git a/clang/test/OpenMP/target_update_strided_messages.c b/clang/test/OpenMP/target_update_strided_messages_from.c
similarity index 92%
rename from clang/test/OpenMP/target_update_strided_messages.c
rename to clang/test/OpenMP/target_update_strided_messages_from.c
index 1f50af4e52805..7696505cc58f1 100644
--- a/clang/test/OpenMP/target_update_strided_messages.c
+++ b/clang/test/OpenMP/target_update_strided_messages_from.c
@@ -1,8 +1,6 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-void foo(void) {}
-
int main(int argc, char **argv) {
int len = 8;
double data[len];
@@ -11,9 +9,6 @@ int main(int argc, char **argv) {
#pragma omp target update from(data[0:4:2]) // OK
{}
- #pragma omp target update to(data[0:len/2:2]) // OK
- {}
-
#pragma omp target update from(data[1:3:2]) // OK
{}
diff --git a/clang/test/OpenMP/target_update_strided_messages_to.c b/clang/test/OpenMP/target_update_strided_messages_to.c
new file mode 100644
index 0000000000000..e765c6344d1b5
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_messages_to.c
@@ -0,0 +1,33 @@
+// 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 = 8;
+ double data[len];
+
+ // Valid strided array sections
+ #pragma omp target update to(data[0:4:2]) // OK
+ {}
+
+ #pragma omp target update to(data[1:3:2]) // OK
+ {}
+
+ // Missing stride (default = 1)
+ #pragma omp target update to(data[0:4]) // OK
+ {}
+
+ // Invalid stride expressions
+ #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'}}
+
+ #pragma omp target update to(data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Missing colon
+ #pragma omp target update to(data[0:4 2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+
+ // Too many colons
+ #pragma omp target update to(data[0:4:2:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_multiple_messages.c b/clang/test/OpenMP/target_update_strided_multiple_messages_from.c
similarity index 60%
rename from clang/test/OpenMP/target_update_strided_multiple_messages.c
rename to clang/test/OpenMP/target_update_strided_multiple_messages_from.c
index 361d4c66c362b..2cb62007c4ea4 100644
--- a/clang/test/OpenMP/target_update_strided_multiple_messages.c
+++ b/clang/test/OpenMP/target_update_strided_multiple_messages_from.c
@@ -1,44 +1,21 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-void foo(void) {}
-
-typedef struct {
- int len;
- double data[12];
-} S;
-
int main(int argc, char **argv) {
int len = 12;
double data1[len], data2[len];
- S s;
// Valid multiple strided array sections
#pragma omp target update from(data1[0:4:2], data2[0:2:5]) // OK
{}
- #pragma omp target update to(data1[1:2:3], data2[2:3:2]) // OK
- {}
-
// Mixed strided and regular array sections
#pragma omp target update from(data1[0:len], data2[0:4:2]) // OK
{}
- // Struct member arrays with strides
- #pragma omp target update from(s.data[0:4:2]) // OK
- {}
-
- #pragma omp target update from(s.data[0:s.len/2:2]) // OK
- {}
-
// Invalid stride in one of multiple sections
#pragma omp target update from(data1[0:3:4], data2[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
- // Complex expressions in multiple arrays
- int stride1 = 2, stride2 = 3;
- #pragma omp target update from(data1[0:len/2:stride1], data2[1:len/3:stride2]) // OK
- {}
-
// Missing colon
#pragma omp target update from(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
diff --git a/clang/test/OpenMP/target_update_strided_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_multiple_messages_to.c
new file mode 100644
index 0000000000000..378da16286c24
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_multiple_messages_to.c
@@ -0,0 +1,25 @@
+// 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 = 12;
+ double data1[len], data2[len];
+
+ // Valid multiple strided array sections
+ #pragma omp target update to(data1[0:4:2], data2[0:2:5]) // OK
+ {}
+
+ // Mixed strided and regular array sections
+ #pragma omp target update to(data1[0:len], data2[0:4:2]) // OK
+ {}
+
+ // Invalid stride in one of multiple sections
+ #pragma omp target update to(data1[0:3:4], data2[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+ {}
+
+ // Missing colon
+ #pragma omp target update to(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+ {}
+
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_partial_messages.c b/clang/test/OpenMP/target_update_strided_partial_messages_from.c
similarity index 59%
rename from clang/test/OpenMP/target_update_strided_partial_messages.c
rename to clang/test/OpenMP/target_update_strided_partial_messages_from.c
index 6dc286c8a1161..faacb99ae2404 100644
--- a/clang/test/OpenMP/target_update_strided_partial_messages.c
+++ b/clang/test/OpenMP/target_update_strided_partial_messages_from.c
@@ -1,7 +1,6 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-
-void foo(void) {}
+// expected-no-diagnostics
int main(int argc, char **argv) {
int len = 11;
@@ -17,16 +16,10 @@ int main(int argc, char **argv) {
// Valid: complex expressions
int offset = 1;
- int count = 3;
- int stride = 2;
- #pragma omp target update from(data[offset:count:stride]) // OK
- {}
-
+
// Invalid stride expressions
#pragma omp target update from(data[0:4:offset-1]) // OK if offset > 1
{}
- #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'}}
-
return 0;
}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_partial_messages_to.c
new file mode 100644
index 0000000000000..9d02269bd610e
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_partial_messages_to.c
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+// expected-no-diagnostics
+
+int main(int argc, char **argv) {
+ int len = 11;
+ double data[len];
+
+ // Valid partial strided updates
+ #pragma omp target update to(data[0:4:3]) // OK
+ {}
+
+ // Stride larger than length
+ #pragma omp target update to(data[0:2:10]) // OK
+ {}
+
+ // Valid: complex expressions
+ int offset = 1;
+
+ // Potentially invalid stride expressions depending on runtime values
+ #pragma omp target update to(data[0:4:offset-1]) // OK if offset > 1
+ {}
+
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/strided_multiple_update.c b/offload/test/offloading/strided_multiple_update_from.c
similarity index 100%
rename from offload/test/offloading/strided_multiple_update.c
rename to offload/test/offloading/strided_multiple_update_from.c
diff --git a/offload/test/offloading/strided_multiple_update_to.c b/offload/test/offloading/strided_multiple_update_to.c
new file mode 100644
index 0000000000000..cfe4513514fa9
--- /dev/null
+++ b/offload/test/offloading/strided_multiple_update_to.c
@@ -0,0 +1,122 @@
+// This test checks that #pragma omp target update to(data1[0:3:4],
+// data2[0:2:5]) correctly updates disjoint strided sections of multiple arrays
+// from the host to the device.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int len = 12;
+ double data1[len], data2[len];
+
+ // Initialize host arrays
+ for (int i = 0; i < len; i++) {
+ data1[i] = i;
+ data2[i] = i * 10;
+ }
+
+ printf("original host array values:\n");
+ printf("data1:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data1[i]);
+ printf("data2:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data2[i]);
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: data1:
+ // 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: data2:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 30.0
+ // CHECK-NEXT: 40.0
+ // CHECK-NEXT: 50.0
+ // CHECK-NEXT: 60.0
+ // CHECK-NEXT: 70.0
+ // CHECK-NEXT: 80.0
+ // CHECK-NEXT: 90.0
+ // CHECK-NEXT: 100.0
+ // CHECK-NEXT: 110.0
+
+#pragma omp target data map(tofrom : data1[0 : len], data2[0 : len])
+ {
+ // Initialize device arrays to 20
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ data1[i] = 20.0;
+ data2[i] = 20.0;
+ }
+ }
+
+ // Modify host arrays for strided elements
+ data1[0] = 10.0;
+ data1[4] = 10.0;
+ data1[8] = 10.0;
+ data2[0] = 10.0;
+ data2[5] = 10.0;
+
+ // data1[0:3:4] // indices 0,4,8
+ // data2[0:2:5] // indices 0,5
+#pragma omp target update to(data1[0 : 3 : 4], data2[0 : 2 : 5])
+
+ // Verify on device by adding 5
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++)
+ data1[i] += 5.0;
+ for (int i = 0; i < len; i++)
+ data2[i] += 5.0;
+ }
+ }
+
+ printf("device array values after update to:\n");
+ printf("data1:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data1[i]);
+ printf("data2:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data2[i]);
+
+ // CHECK: device array values after update to:
+ // CHECK-NEXT: data1:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: data2:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+}
diff --git a/offload/test/offloading/strided_partial_update.c b/offload/test/offloading/strided_partial_update_from.c
similarity index 100%
rename from offload/test/offloading/strided_partial_update.c
rename to offload/test/offloading/strided_partial_update_from.c
diff --git a/offload/test/offloading/strided_partial_update_to.c b/offload/test/offloading/strided_partial_update_to.c
new file mode 100644
index 0000000000000..e03f3a9f9f607
--- /dev/null
+++ b/offload/test/offloading/strided_partial_update_to.c
@@ -0,0 +1,72 @@
+// This test checks that #pragma omp target update to(data[0:4:3]) correctly
+// updates every third element (stride 3) from the host to the device, partially
+// across the array
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int len = 11;
+ 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");
+
+ // 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
+
+#pragma omp target data map(tofrom : data[0 : len])
+ {
+ // Initialize device array to 20
+#pragma omp target
+ for (int i = 0; i < len; i++)
+ data[i] = 20.0;
+
+ // Modify host data for strided elements
+ data[0] = 10.0;
+ data[3] = 10.0;
+ data[6] = 10.0;
+ data[9] = 10.0;
+
+#pragma omp target update to(data[0 : 4 : 3]) // indices 0,3,6,9
+
+ // Verify on device by adding 5
+#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]);
+ printf("\n");
+
+ // CHECK: 15.000000
+ // CHECK: 25.000000
+ // CHECK: 25.000000
+ // CHECK: 15.000000
+ // CHECK: 25.000000
+ // CHECK: 25.000000
+ // CHECK: 15.000000
+ // CHECK: 25.000000
+ // CHECK: 25.000000
+ // CHECK: 15.000000
+ // CHECK: 25.000000
+}
diff --git a/offload/test/offloading/strided_update.c b/offload/test/offloading/strided_update_from.c
similarity index 100%
rename from offload/test/offloading/strided_update.c
rename to offload/test/offloading/strided_update_from.c
diff --git a/offload/test/offloading/strided_update_to.c b/offload/test/offloading/strided_update_to.c
new file mode 100644
index 0000000000000..89f80e021394a
--- /dev/null
+++ b/offload/test/offloading/strided_update_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:4: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 : 4 : 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;
+}
\ No newline at end of file
More information about the llvm-commits
mailing list