[clang] [llvm] [Clang][OpenMP] Support expression semantics in target update fields with non-contiguous array sections (PR #176708)

[Amit Tiwari via cfe-commits]

================
@@ -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):
----------------
amitamd7 wrote:

For now, I have opted for the first option of separating them out with a function.

https://github.com/llvm/llvm-project/pull/176708