[Offload][Tests] Non-contiguous_update_to_tests #169623

amitamd7 · 2025-11-26T09:31:28Z

PR #144635 enabled non-contiguous updates for both update from and update to clauses, but tests for update to were missing. This PR adds those missing tests to ensure coverage.

github-actions · 2025-11-26T09:33:24Z

✅ With the latest revision this PR passed the C/C++ code formatter.

llvmbot · 2025-11-26T15:52:21Z

@llvm/pr-subscribers-offload

Author: Amit Tiwari (amitamd7)

Changes

PR #144635 enabled non-contiguous updates for both update from and update to clauses, but tests for update to were missing. This PR adds those missing tests to ensure coverage.

Full diff: https://github.com/llvm/llvm-project/pull/169623.diff

6 Files Affected:

(renamed) offload/test/offloading/strided_multiple_update_from.c ()
(added) offload/test/offloading/strided_multiple_update_to.c (+122)
(renamed) offload/test/offloading/strided_partial_update_from.c ()
(added) offload/test/offloading/strided_partial_update_to.c (+72)
(renamed) offload/test/offloading/strided_update_from.c ()
(added) offload/test/offloading/strided_update_to.c (+72)

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..1ae61e9ed006c
--- /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
+}
\ No newline at end of file
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..ea708cf0c79c4
--- /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
+}
\ No newline at end of file
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

Copilot

Pull request overview

This PR adds test coverage for non-contiguous update to clauses in OpenMP offloading, complementing PR #144635 which enabled this functionality. The tests verify that strided array updates correctly transfer data from host to device with various stride patterns.

Key changes:

Added tests for basic strided updates with stride-2 pattern
Added tests for partial strided updates with stride-3 pattern
Added tests for multiple arrays with different stride patterns

Reviewed changes

Copilot reviewed 3 out of 6 changed files in this pull request and generated 2 comments.

File	Description
offload/test/offloading/strided_update_to.c	Tests basic strided update with stride-2 (every other element)
offload/test/offloading/strided_partial_update_to.c	Tests partial strided update with stride-3 (every third element)
offload/test/offloading/strided_multiple_update_to.c	Tests multiple arrays with different strides (stride-4 and stride-5)

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Copilot · 2025-11-26T15:57:57Z

offload/test/offloading/strided_update_to.c

+  // CHECK: 0.000000
+  // CHECK: 1.000000
+  // CHECK: 2.000000
+  // CHECK: 3.000000
+  // CHECK: 4.000000
+  // CHECK: 5.000000
+  // CHECK: 6.000000
+  // CHECK: 7.000000


These CHECK comments verify the wrong output. They appear before the 'device array values after update to:' print statement (line 57), but they're checking for the original host values (0-7) instead of the expected device values after the update. These checks will fail because they don't match what's actually printed at this point in the execution.

Copilot · 2025-11-26T15:57:58Z

offload/test/offloading/strided_partial_update_to.c

+  // CHECK: 8.000000
+  // CHECK: 9.000000
+  // CHECK: 10.000000
+


These CHECK comments are positioned after the 'original host array values:' print statement (line 18), but the printf on line 21 includes a newline ('\n') after printing all values. The CHECK directives should account for this empty line or be repositioned to avoid potential matching issues with the test framework.

Suggested change

// CHECK:

saiislam

Can these multiple and highly related patterns for single, partial, and multiple updates be tested in the same file?

saiislam · 2025-11-26T15:59:49Z

offload/test/offloading/strided_multiple_update_to.c

+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+}


Add a last line

code-formatter fails with this change. Hence, reverting it.

saiislam · 2025-11-26T16:00:05Z

offload/test/offloading/strided_partial_update_to.c

+  // CHECK: 25.000000
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+}


Add last line please

amitamd7 · 2025-11-27T08:31:08Z

Can these multiple and highly related patterns for single, partial, and multiple updates be tested in the same file?

No, as these update affect Clang’s internal data structures differently, not only in terms of handling strided data but also in how they manage the number of components involved. Each takes a distinct path in clang, hence, we would want them to work/fail independently.

shiltian · 2025-12-01T16:38:44Z

clang/test/OpenMP/target_update_strided_messages_to.c

+  {}
+
+  return 0;
+}


missing an empty file at EoF

shiltian · 2025-12-01T16:38:52Z

clang/test/OpenMP/target_update_strided_multiple_messages_to.c

+  {}
+
+  return 0;
+}


here as well

shiltian · 2025-12-01T16:38:57Z

clang/test/OpenMP/target_update_strided_partial_messages_from.c

-  #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;
 }


shiltian · 2025-12-01T16:39:12Z

clang/test/OpenMP/target_update_strided_partial_messages_to.c

+  {}
+
+  return 0;
+}


all new files

amitamd7 force-pushed the tiwari_stride_update_tests branch from 6ab9518 to 3baf914 Compare November 26, 2025 09:42

amitamd7 marked this pull request as ready for review November 26, 2025 15:51

llvmbot added the offload label Nov 26, 2025

saiislam requested review from Copilot and jtb20 November 26, 2025 15:57

Copilot AI reviewed Nov 26, 2025

View reviewed changes

saiislam reviewed Nov 26, 2025

View reviewed changes

amitamd7 force-pushed the tiwari_stride_update_tests branch 4 times, most recently from 9aca391 to ab3cced Compare November 27, 2025 08:49

amitamd7 requested a review from saiislam December 1, 2025 06:04

amitamd7 marked this pull request as draft December 1, 2025 06:27

amitamd7 force-pushed the tiwari_stride_update_tests branch from ab3cced to e38e0fe Compare December 1, 2025 08:20

amitamd7 marked this pull request as ready for review December 1, 2025 09:11

llvmbot added the clang:openmp OpenMP related changes to Clang label Dec 1, 2025

shiltian reviewed Dec 1, 2025

View reviewed changes

amitamd7 force-pushed the tiwari_stride_update_tests branch from e38e0fe to d740866 Compare December 1, 2025 17:50

amitamd7 marked this pull request as draft December 1, 2025 17:54

non-contiguous_update_to_tests

4ee4ce4

amitamd7 force-pushed the tiwari_stride_update_tests branch from d740866 to 4ee4ce4 Compare December 1, 2025 19:12

amitamd7 marked this pull request as ready for review December 1, 2025 19:17

[Offload][Tests] Non-contiguous_update_to_tests #169623

Are you sure you want to change the base?

[Offload][Tests] Non-contiguous_update_to_tests #169623

Conversation

amitamd7 commented Nov 26, 2025

Uh oh!

github-actions bot commented Nov 26, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Nov 26, 2025

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Copilot AI Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

Copilot AI Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

saiislam left a comment

Choose a reason for hiding this comment

Uh oh!

saiislam Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

amitamd7 Nov 27, 2025

Choose a reason for hiding this comment

Uh oh!

saiislam Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

amitamd7 commented Nov 27, 2025

Uh oh!

shiltian Dec 1, 2025

Choose a reason for hiding this comment

Uh oh!

shiltian Dec 1, 2025

Choose a reason for hiding this comment

Uh oh!

shiltian Dec 1, 2025

Choose a reason for hiding this comment

Uh oh!

shiltian Dec 1, 2025

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

github-actions bot commented Nov 26, 2025 •

edited

Loading