[Clang][OpenMP] Support expression semantics in target update fields with non-contiguous array sections
#175505
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
clang:codegen
Member
|
@llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-flang-openmp Author: Amit Tiwari (amitamd7) ChangesIssue 1: Dimension override bug with variable count expressions When variable count expressions were used with stride, the constant subsection path computed size first. This marked ArgSizes with byte size semantics. Variable expression logic later triggered, but reused ArgSizes assuming "bytes" semantics Result: ArgSizes wasn't overwritten with dimension count, breaking non-contiguous mapping. Issue 2: Variable stride not recognized as non-contiguous Issue 3: Missing expression semantics in OMPIRBuilder Fixes:
Added testcases to cover stack arrays, heap pointers, struct members, etc for expression semantics in non-contiguous update. Patch is 95.08 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/175505.diff 32 Files Affected:
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index b8ee701c482bb..766a9d24409b3 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7908,11 +7908,27 @@ class MappableExprsHandler {
if (!StrideExpr)
return false;
+ assert(StrideExpr->getType()->isIntegerType() &&
+ "Stride expression must be of integer type");
+
+ // If the stride is a variable (not a constant), it's non-contiguous.
+ const Expr *S = StrideExpr->IgnoreParenImpCasts();
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(S)) {
+ if (isa<VarDecl>(DRE->getDecl()) ||
+ isa<ParmVarDecl>(DRE->getDecl()))
+ return true;
+ }
+ if (isa<MemberExpr>(S) || isa<ArraySubscriptExpr>(S))
+ return true;
+
+ // If stride is not evaluatable as a constant, treat as
+ // non-contiguous.
const auto Constant =
StrideExpr->getIntegerConstantExpr(CGF.getContext());
if (!Constant)
return false;
+ // Treat non-unitary strides as non-contiguous.
return !Constant->isOne();
});
diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c
new file mode 100644
index 0000000000000..932cd6b1c97bb
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c
@@ -0,0 +1,62 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+int main(int argc, char **argv) {
+ int len = 16;
+ int count = 8;
+ int stride = 2;
+ int stride_large = 5;
+ double *data;
+
+ // Valid strided array sections with both variable count and variable stride (FROM)
+ #pragma omp target update from(data[0:count:stride]) // OK - both variable
+ {}
+
+ #pragma omp target update from(data[0:len/2:stride]) // OK - count expression, variable stride
+ {}
+
+ #pragma omp target update from(data[0:count:stride_large]) // OK - variable count, different stride
+ {}
+
+ #pragma omp target update from(data[1:len-2:stride]) // OK - with offset, count expression
+ {}
+
+ #pragma omp target update from(data[0:count/2:stride*2]) // OK - both expressions
+ {}
+
+ #pragma omp target update from(data[0:(len+1)/2:stride+1]) // OK - complex expressions
+ {}
+
+ #pragma omp target update from(data[2:count-2:len/4]) // OK - all expressions
+ {}
+
+ // Edge cases
+ int stride_one = 1;
+ #pragma omp target update from(data[0:count:stride_one]) // OK - variable count, stride=1
+ {}
+
+ #pragma omp target update from(data[0:len/stride:stride]) // OK - count depends on stride
+ {}
+
+ // Invalid compile-time constant strides with variable count
+ #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'}}
+
+ #pragma omp target update from(data[0:len/2:-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'}}
+
+ #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Valid strided array sections with variable count and stride (TO)
+ #pragma omp target update to(data[0:count:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/2:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:count:stride*2]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(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;
+}
diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c
new file mode 100644
index 0000000000000..23fba9c8bc84f
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+int main(int argc, char **argv) {
+ int len = 16;
+ int count = 8;
+ int divisor = 2;
+ double *data;
+
+ // Valid strided array sections with variable count expressions (FROM)
+ #pragma omp target update from(data[0:count:2]) // OK - variable count
+ {}
+
+ #pragma omp target update from(data[0:len/2:2]) // OK - count expression
+ {}
+
+ #pragma omp target update from(data[0:len-4:3]) // OK - count with subtraction
+ {}
+
+ #pragma omp target update from(data[1:(len+1)/2:2]) // OK - complex count expression
+ {}
+
+ #pragma omp target update from(data[0:count*2:3]) // OK - count multiplication
+ {}
+
+ #pragma omp target update from(data[2:len%divisor:2]) // OK - count with modulo
+ {}
+
+ // Variable count with stride = 1 (contiguous)
+ #pragma omp target update from(data[0:count]) // OK - variable count, implicit stride
+ {}
+
+ #pragma omp target update from(data[0:len/divisor]) // OK - expression count, implicit stride
+ {}
+
+ // Invalid stride expressions with variable count
+ #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'}}
+
+ #pragma omp target update from(data[0:len/2:-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'}}
+
+ #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Valid strided array sections with variable count expressions (TO)
+ #pragma omp target update to(data[0:count:2]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/2:2]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len-4:3]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(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;
+}
diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c
new file mode 100644
index 0000000000000..3f85ed0c48d66
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c
@@ -0,0 +1,64 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+int main(int argc, char **argv) {
+ int len = 16;
+ int stride = 2;
+ int stride_large = 5;
+ double *data;
+
+ // Valid strided array sections with variable stride (FROM)
+ #pragma omp target update from(data[0:8:stride]) // OK - variable stride
+ {}
+
+ #pragma omp target update from(data[0:4:stride_large]) // OK - different variable stride
+ {}
+
+ #pragma omp target update from(data[1:6:stride]) // OK - with offset
+ {}
+
+ #pragma omp target update from(data[0:5:stride+1]) // OK - stride expression
+ {}
+
+ #pragma omp target update from(data[0:4:stride*2]) // OK - stride multiplication
+ {}
+
+ #pragma omp target update from(data[2:3:len/4]) // OK - stride from expression
+ {}
+
+ // Edge case: stride = 1 (should be contiguous, not non-contiguous)
+ int stride_one = 1;
+ #pragma omp target update from(data[0:8:stride_one]) // OK - stride=1 is contiguous
+ {}
+
+ // Invalid variable stride expressions
+ int zero_stride = 0;
+ int neg_stride = -1;
+
+ // Note: These are runtime checks, so no compile-time error
+ #pragma omp target update from(data[0:8:zero_stride]) // OK at compile-time (runtime will fail)
+ {}
+
+ #pragma omp target update from(data[0:4:neg_stride]) // OK at compile-time (runtime will fail)
+ {}
+
+ // Compile-time constant invalid strides
+ #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'}}
+
+ // Valid strided array sections with variable stride (TO)
+ #pragma omp target update to(data[0:8:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:5:stride+1]) // OK
+ {}
+
+ #pragma omp target update to(data[0:4:stride*2]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(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'}}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c
new file mode 100644
index 0000000000000..70775d5c8322c
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c
@@ -0,0 +1,72 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+#define N 20
+typedef struct {
+ double data[N];
+ int len;
+ int stride;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = 16;
+ s.stride = 2;
+ int count = 8;
+ int ext_stride = 3;
+
+ // Valid strided struct member array sections with variable count/stride (FROM)
+ #pragma omp target update from(s.data[0:s.len/2:2]) // OK - member count expression
+ {}
+
+ #pragma omp target update from(s.data[0:count:s.stride]) // OK - external count, member stride
+ {}
+
+ #pragma omp target update from(s.data[0:s.len:ext_stride]) // OK - member count, external stride
+ {}
+
+ #pragma omp target update from(s.data[0:count:ext_stride]) // OK - both external
+ {}
+
+ #pragma omp target update from(s.data[0:s.len/2:s.stride]) // OK - both from struct
+ {}
+
+ #pragma omp target update from(s.data[1:(s.len-2)/2:s.stride]) // OK - complex count expression
+ {}
+
+ #pragma omp target update from(s.data[0:count*2:s.stride+1]) // OK - expressions for both
+ {}
+
+ // Edge cases
+ int stride_one = 1;
+ #pragma omp target update from(s.data[0:s.len:stride_one]) // OK - stride=1
+ {}
+
+ #pragma omp target update from(s.data[0:s.len/s.stride:s.stride]) // OK - count depends on stride
+ {}
+
+ // Invalid compile-time constant strides with variable count
+ #pragma omp target update from(s.data[0:s.len: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(s.data[0:count:-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'}}
+
+ #pragma omp target update from(s.data[1:s.len/2:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Valid strided struct member array sections with variable count and stride (TO)
+ #pragma omp target update to(s.data[0:s.len/2:2]) // OK
+ {}
+
+ #pragma omp target update to(s.data[0:count:s.stride]) // OK
+ {}
+
+ #pragma omp target update to(s.data[0:s.len:ext_stride]) // OK
+ {}
+
+ #pragma omp target update to(s.data[0:count*2:s.stride+1]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(s.data[0:s.len: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;
+}
diff --git a/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c
new file mode 100644
index 0000000000000..0082539538a32
--- /dev/null
+++ b/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c
@@ -0,0 +1,85 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+int main(int argc, char **argv) {
+ int len = 16;
+ int count = 8;
+ int stride = 2;
+ int divisor = 2;
+ double data[100];
+
+ // Valid strided array sections with variable count expressions (FROM)
+ #pragma omp target update from(data[0:count:2]) // OK - variable count
+ {}
+
+ #pragma omp target update from(data[0:len/2:2]) // OK - count expression
+ {}
+
+ #pragma omp target update from(data[0:len-4:3]) // OK - count with subtraction
+ {}
+
+ #pragma omp target update from(data[1:(len+1)/2:2]) // OK - complex count expression
+ {}
+
+ #pragma omp target update from(data[0:count*2:3]) // OK - count multiplication
+ {}
+
+ #pragma omp target update from(data[2:len%divisor:2]) // OK - count with modulo
+ {}
+
+ // Variable stride with constant/variable count
+ #pragma omp target update from(data[0:10:stride]) // OK - constant count, variable stride
+ {}
+
+ #pragma omp target update from(data[0:count:stride]) // OK - both variable
+ {}
+
+ #pragma omp target update from(data[0:len/2:stride]) // OK - count expression, variable stride
+ {}
+
+ #pragma omp target update from(data[0:count:stride*2]) // OK - variable count, stride expression
+ {}
+
+ #pragma omp target update from(data[0:len/divisor:stride+1]) // OK - both expressions
+ {}
+
+ // Variable count with stride = 1 (contiguous)
+ #pragma omp target update from(data[0:count]) // OK - variable count, implicit stride
+ {}
+
+ #pragma omp target update from(data[0:len/divisor]) // OK - expression count, implicit stride
+ {}
+
+ // Edge cases
+ int stride_one = 1;
+ #pragma omp target update from(data[0:len:stride_one]) // OK - stride=1 variable
+ {}
+
+ #pragma omp target update from(data[0:len/stride:stride]) // OK - count depends on stride
+ {}
+
+ // Invalid stride expressions with variable count
+ #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'}}
+
+ #pragma omp target update from(data[0:len/2:-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'}}
+
+ #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Valid strided array sections with variable count expressions (TO)
+ #pragma omp target update to(data[0:count:2]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/2:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:count:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/divisor:stride+1]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(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;
+}
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 716f8582dd7b2..fe6c755b0e504 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9365,16 +9365,29 @@ Error OpenMPIRBuilder::emitOffloadingArrays(
ConstantInt::get(Int64Ty, 0));
SmallBitVector RuntimeSizes(CombinedInfo.Sizes.size());
for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) {
+ bool IsNonContigEntry =
+ IsNonContiguous &&
+ (static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ CombinedInfo.Types[I] &
+ OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0);
+ // For NON_CONTIG entries ArgSizes must carry the dimension count
+ // (number of descriptor_dim records) – NOT the byte size expression.
+ // Variable subsection forms (e.g. 0:s.len/2:2) previously produced a
+ // non-constant size so we marked them runtime and stored the byte size,
+ // leading the runtime to treat it as DimSize and overrun descriptors.
+ if (IsNonContigEntry) {
+ // Dims must be long enough and positive.
+ assert(I < CombinedInfo.NonContigInfo.Dims.size() &&
+ "Induction variable is in-bounds with the NON_CONTIG Dims array");
+ const uint64_t DimCount = CombinedInfo.NonContigInfo.Dims[I];
+ assert(DimCount > 0 && "NON_CONTIG DimCount must be > 0");
+ ConstSizes[I] =
+ ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]);
+ continue;
+ }
if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) {
if (!isa<ConstantExpr>(CI) && !isa<GlobalValue>(CI)) {
- if (IsNonContiguous &&
- static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- CombinedInfo.Types[I] &
- OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG))
- ConstSizes[I] =
- ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]);
- else
- ConstSizes[I] = CI;
+ ConstSizes[I] = CI;
continue;
}
}
diff --git a/offload/test/offloading/strided_update_count_expression.c b/offload/test/offloading/strided_update_count_expression.c
new file mode 100644
index 0000000000000..a87da289a9154
--- /dev/null
+++ b/offload/test/offloading/strided_update_count_expression.c
@@ -0,0 +1,133 @@
+// This test checks that "update from" and "update to" clauses in OpenMP are
+// supported when elements are updated in a non-contiguous manner with variable
+// count expression. Tests #pragma omp target update from/to(data[0:len/2:2])
+// where the count (len/2) is a variable expression, not a constant.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int len = 10;
+ double data[len];
+
+ // ====================================================================
+ // TEST 1: Update FROM device (device -> host)
+ // ====================================================================
+
+#pragma omp target map(tofrom : len, data[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ }
+ }
+
+ printf("Test 1: Update FROM device\n");
+ printf("original host array values:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+
+#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 : len / 2 : 2])
+ }
+
+ printf("from target array results:\n");
+ f...
[truncated]
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
amitamd7
force-pushed
the
tiwari_expression_semantics_non-contig_update
branch
from
db32432 to
04a311b
Compare
amitamd7
force-pushed
the
tiwari_expression_semantics_non-contig_update
branch
from
04a311b to
25ad443
Compare
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Issue 1: Dimension override bug with variable count expressions
When variable count expressions were used with stride, the constant subsection path computed size first. This marked ArgSizes with byte size semantics. Variable expression logic later triggered, but reused ArgSizes assuming "bytes" semantics
Result: ArgSizes wasn't overwritten with dimension count, breaking non-contiguous mapping.
Issue 2: Variable stride not recognized as non-contiguous
CGOpenMPRuntime.cppfailed to detectDeclRefExpr, MemberExpr, ArraySubscriptExpras non-contiguous.Issue 3: Missing expression semantics in OMPIRBuilder
OMPIRBuilder.cppdidn't handle dimension count forOMP_MAP_NON_CONTIGflagFixes:
clang/lib/CodeGen/CGOpenMPRuntime.cpp- Variable stride detection + dimension count logicllvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp- Expression semantics for non-contiguousDeclRefExpr/MemberExpr/ArraySubscriptExpr) as non-contiguousOMP_MAP_NON_CONTIGflag (0x100000000000) for variable stride/count.Added testcases to cover stack arrays, heap pointers, struct members, etc., for expression semantics in non-contiguous update.