[Clang][OpenMP] Non-contiguous strided update

clang/docs/OpenMPSupport.rst

@@ -191,7 +191,7 @@ implementation.
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | device                       | teams construct on the host device                           | :good:`done`             | r371553                                                               |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
-| device                       | support non-contiguous array sections for target update      | :good:`done`             |                                                                       |
+| device                       | support non-contiguous array sections for target update      | :good:`done`             | https://github.com/llvm/llvm-project/pull/144635                                                                      |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | device                       | pointer attachment                                           | :good:`done`             |                                                                       |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+

clang/docs/ReleaseNotes.rst

@@ -262,6 +262,7 @@ OpenMP Support
 - Added parsing and semantic analysis support for the ``need_device_addr``
   modifier in the ``adjust_args`` clause.
 - Allow array length to be omitted in array section subscript expression.
+- Fixed non-contiguous strided update in the ``omp target update`` directive with the ``from`` clause.
 
 Improvements
 ^^^^^^^^^^^^

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -7487,7 +7487,32 @@ class MappableExprsHandler {
     // dimension.
     uint64_t DimSize = 1;
 
-    bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
+    // Detects non-contiguous updates due to strided accesses.
+    // Sets the 'IsNonContiguous' flag so that the 'MapType' bits are set
+    // correctly when generating information to be passed to the runtime. The
+    // flag is set to true if any array section has a stride not equal to 1, or
+    // if the stride is not a constant expression (conservatively assumed
+    // non-contiguous).
+    bool IsNonContiguous =
+        CombinedInfo.NonContigInfo.IsNonContiguous ||
+        any_of(Components, [&](const auto &Component) {
+          const auto *OASE =
+              dyn_cast<ArraySectionExpr>(Component.getAssociatedExpression());
+          if (!OASE)
+            return false;
+
+          const Expr *StrideExpr = OASE->getStride();
+          if (!StrideExpr)
+            return false;
+
+          const auto Constant =
+              StrideExpr->getIntegerConstantExpr(CGF.getContext());
+          if (!Constant)
+            return false;
+
+          return !Constant->isOne();
+        });
+
     bool IsPrevMemberReference = false;
 
     bool IsPartialMapped =

clang/test/OpenMP/target_update_strided_messages.c

@@ -0,0 +1,38 @@
+// 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];
+
+  // Valid strided array sections
+  #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
+  {}
+
+  // Missing stride (default = 1)
+  #pragma omp target update from(data[0:4]) // OK
+  {}
+
+  // Invalid stride expressions
+  #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'}}
+
+  // Missing colon 
+  #pragma omp target update from(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 from(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;
+}

clang/test/OpenMP/target_update_strided_multiple_messages.c

@@ -0,0 +1,46 @@
+// 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 '['}}
+
+  return 0;
+}

clang/test/OpenMP/target_update_strided_partial_messages.c

@@ -0,0 +1,32 @@
+// 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 = 11;
+  double data[len];
+
+  // Valid partial strided updates
+  #pragma omp target update from(data[0:4:3]) // OK
+  {}
+
+  // Stride larger than length
+  #pragma omp target update from(data[0:2:10]) // OK
+  {}
+
+  // 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;
+}

offload/test/offloading/strided_multiple_update.c

@@ -0,0 +1,62 @@
+// This test checks that #pragma omp target update from(data1[0:3:4],
+// data2[0:2:5]) correctly updates disjoint strided sections of multiple arrays
+// from the device to the host.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int len = 12;
+  double data1[len], data2[len];
+
+// Initial values
+#pragma omp target map(tofrom : data1[0 : len], data2[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      data1[i] = i;
+      data2[i] = i * 10;
+    }
+  }
+
+  printf("original host array values:\n");
+  printf("data1: ");
+  for (int i = 0; i < len; i++)
+    printf("%.1f ", data1[i]);
+  printf("\ndata2: ");
+  for (int i = 0; i < len; i++)
+    printf("%.1f ", data2[i]);
+  printf("\n\n");
+
+#pragma omp target data map(to : data1[0 : len], data2[0 : len])
+  {
+// Modify arrays on device
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++)
+        data1[i] += i;
+      for (int i = 0; i < len; i++)
+        data2[i] += 100;
+    }
+
+// data1[0:3:4]  // indices 0,4,8
+// data2[0:2:5]  // indices 0,5
+#pragma omp target update from(data1[0 : 3 : 4], data2[0 : 2 : 5])
+  }
+
+  printf("device array values after update from:\n");
+  printf("data1: ");
+  for (int i = 0; i < len; i++)
+    printf("%.1f ", data1[i]);
+  printf("\ndata2: ");
+  for (int i = 0; i < len; i++)
+    printf("%.1f ", data2[i]);
+  printf("\n\n");
+
+  // CHECK: data1: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0
+  // CHECK: data2: 0.0 10.0 20.0 30.0 40.0 50.0 60.0 70.0 80.0 90.0 100.0 110.0
+
+  // CHECK: data1: 0.0 1.0 2.0 3.0 8.0 5.0 6.0 7.0 16.0 9.0 10.0 11.0
+  // CHECK: data2: 100.0 10.0 20.0 30.0 40.0 150.0 60.0 70.0 80.0 90.0 100.0
+  // 110.0
+}

offload/test/offloading/strided_partial_update.c

@@ -0,0 +1,63 @@
+// This test checks that #pragma omp target update from(data[0:4:3]) correctly
+// updates every third element (stride 3) from the device to the host, 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];
+
+#pragma omp target map(tofrom : 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 : 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 : 4 : 3]) // indices 0,3,6,9
+  }
+
+  printf("device array values after update from:\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
+
+  // CHECK: 0.000000
+  // CHECK: 1.000000
+  // CHECK: 2.000000
+  // CHECK: 6.000000
+  // CHECK: 4.000000
+  // CHECK: 5.000000
+  // CHECK: 12.000000
+  // CHECK: 7.000000
+  // CHECK: 8.000000
+  // CHECK: 18.000000
+  // CHECK: 10.000000
+}

offload/test/offloading/strided_update.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:4: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 : 4 : 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;
+}