[Clang][OpenMP] Non-contiguous strided update

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -7384,7 +7384,40 @@ 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 = false;
+    for (const auto &Component : Components) {
+      const auto *OASE =
+          dyn_cast<ArraySectionExpr>(Component.getAssociatedExpression());
+      if (OASE) {
+        const Expr *StrideExpr = OASE->getStride();
+        if (StrideExpr) {
+          // Check if the stride is a constant integer expression
+          if (StrideExpr->isIntegerConstantExpr(CGF.getContext())) {
+            if (auto Constant =
+                    StrideExpr->getIntegerConstantExpr(CGF.getContext())) {
+              int64_t StrideVal = Constant->getExtValue();
+              if (StrideVal != 1) {
+                // Set flag if stride is not 1 (i.e., non-contiguous update)
+                IsNonContiguous = true;
+                break;
+              }
+            }
+          } else {
+            // If stride is not a constant, conservatively treat as
+            // non-contiguous
+            IsNonContiguous = true;
+            break;
+          }
+        }
+      }
+    }
+
     bool IsPrevMemberReference = false;
 
     bool IsPartialMapped =

offload/test/offloading/strided_update.c

@@ -0,0 +1,51 @@
+// Checks that "update from" clause in OpenMP is supported when the elements are updated in a non-contiguous manner.
+// 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])  
+  {  
+    #pragma omp target  
+    for (int i = 0; i < len; i++) {  
+      data[i] += i ;  
+    }  
+
+    #pragma omp target update from(data[0:8:2])  
+  }  
+  // from results  
+  // 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;  
+}  
+