[Clang][OpenMP] Non-contiguous strided update

[amitamd7]

This patch handles the strided update in the #pragma omp target update from(data[a:b:c]) directive where 'c' represents the strided access leading to non-contiguous update in the data array when the offloaded execution returns the control back to host from device using the from clause.

Issue: Clang CodeGen where info is generated for the particular MapType (to, from, etc), it was failing to detect the strided access. Because of this, the MapType bits were incorrect when passed to runtime. This led to incorrect execution (contiguous) in the libomptarget runtime code.

Added a minimal testcase that verifies the working of the patch.

[github-actions]

Thank you for submitting a Pull Request (PR) to the LLVM Project!

This PR will be automatically labeled and the relevant teams will be notified.

If you wish to, you can add reviewers by using the "Reviewers" section on this page.

If this is not working for you, it is probably because you do not have write permissions for the repository. In which case you can instead tag reviewers by name in a comment by using @ followed by their GitHub username.

If you have received no comments on your PR for a week, you can request a review by "ping"ing the PR by adding a comment “Ping”. The common courtesy "ping" rate is once a week. Please remember that you are asking for valuable time from other developers.

If you have further questions, they may be answered by the LLVM GitHub User Guide.

You can also ask questions in a comment on this PR, on the LLVM Discord or on the forums.

[llvmbot]

@llvm/pr-subscribers-offload
@llvm/pr-subscribers-clang-codegen

@llvm/pr-subscribers-clang

Author: Amit Tiwari (amitamd7)

Changes

This patch handles the strided update in the #pragma omp target update from(data[a:b:c]) directive where 'c' represents the strided access leading to non-contiguous update in the data array when the offloaded execution returns the control back to host from device using the from clause.

Issue: Clang CodeGen where info is generated for the particular MapType (to, from, etc), it was failing to detect the strided access. Because of this, the MapType bits were incorrect when passed to runtime. This led to incorrect execution (contiguous) in the libomptarget runtime code.

Added a minimal testcase that verifies the working of the patch.

---

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

2 Files Affected:

• (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+34-1)
• (added) offload/test/offloading/strided_update.c (+51)

diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 4173355491fd4..81a2dd0fae5c9 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/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 =
diff --git a/offload/test/offloading/strided_update.c b/offload/test/offloading/strided_update.c
new file mode 100644
index 0000000000000..fc47216fb5684
--- /dev/null
+++ b/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;  
+}  
+