[Clang][OpenMP] Handle check for non-contiguous mapping in pointer-based array sections

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -8239,13 +8239,25 @@ class MappableExprsHandler {
           ElementType = CAT->getElementType().getTypePtr();
         else if (VAT)
           ElementType = VAT->getElementType().getTypePtr();
-        else
-          assert(&Component == &*Components.begin() &&
-                 "Only expect pointer (non CAT or VAT) when this is the "
-                 "first Component");
-        // If ElementType is null, then it means the base is a pointer
-        // (neither CAT nor VAT) and we'll attempt to get ElementType again
-        // for next iteration.
+        else if (&Component == &*Components.begin()) {
+          // If the base is a raw pointer (e.g. T *data with data[a:b:c]),
+          // there was no earlier CAT/VAT/array handling to establish
+          // ElementType. Capture the pointee type now so that subsequent
+          // components (offset/length/stride) have a concrete element type to
+          // work with. This makes pointer-backed sections behave consistently
+          // with CAT/VAT/array bases.
+          if (const auto *PtrType = Ty->getAs<PointerType>())
+            ElementType = PtrType->getPointeeType().getTypePtr();
+        } else {
+          // Any component after the first should never have a raw pointer type;
+          // by this point. ElementType must already be known (set above or in
+          // prior array / CAT / VAT handling).
+          assert(!Ty->isPointerType() &&
+                 "Non-first components should not be raw pointers");
+        }
+
+        // At this stage, if ElementType was a base pointer and we are in the
+        // first iteration, it has been computed.
         if (ElementType) {
           // For the case that having pointer as base, we need to remove one
           // level of indirection.
@@ -8955,8 +8967,15 @@ class MappableExprsHandler {
       // If there is an entry in PartialStruct it means we have a struct with
       // individual members mapped. Emit an extra combined entry.
       if (PartialStruct.Base.isValid()) {
-        UnionCurInfo.NonContigInfo.Dims.push_back(0);
-        // Emit a combined entry:
+        // Prepend a synthetic dimension of length 1 to represent the
+        // aggregated struct object. Using 1 (not 0, as 0 produced an
+        //    incorrect non-contiguous descriptor (DimSize==1), causing the
+        //    non-contiguous motion clause path to be skipped.) is important:
+        //  * It preserves the correct rank so targetDataUpdate() computes
+        //    DimSize == 2 for cases like strided array sections originating
+        //    from user-defined mappers (e.g. test with s.data[0:8:2]).
+        UnionCurInfo.NonContigInfo.Dims.insert(
+            UnionCurInfo.NonContigInfo.Dims.begin(), 1);
         emitCombinedEntry(CombinedInfo, UnionCurInfo.Types, PartialStruct,
                           /*IsMapThis*/ !VD, OMPBuilder, VD);
       }

clang/test/OpenMP/target_update_codegen.cpp

@@ -1177,7 +1177,7 @@ void foo(int arg) {
 // CK21: [[STRUCT_ST:%.+]] = type { [10 x [10 x [10 x ptr]]] }
 // CK21: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
 
-// CK21: [[SIZE:@.+]] = private unnamed_addr constant [2 x i64] zeroinitializer
+// CK21: [[SIZE:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]
 // CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 299067162755073]
 
 struct ST {
@@ -1221,11 +1221,11 @@ struct ST {
     // CK21: store i64 1, ptr [[COUNT_4]],
     // CK21: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
     // CK21: store i64 {{4|8}}, ptr [[STRIDE_4]],
-    // CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPSZ:%.+]], ptr [[MTYPE]]{{.+}})
-    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
-    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-    // CK21-DAG: [[PTRS:%.+]] = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0
-    // CK21-DAG: store ptr [[DIMS]], ptr [[PTRS]],
+    // CK21-DAG: [[GEPBP:%.+]] = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
+    // CK21-DAG: [[GEPP:%.+]]  = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0
+    // CK21-DAG: [[GEPSZ:%.+]] = getelementptr inbounds [2 x i64], ptr %.offload_sizes, i32 0, i32 0
+    // CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP]], ptr [[GEPP]], ptr [[GEPSZ]], ptr @.offload_maptypes, ptr null, ptr null)
+    // CK21: ret void
 #pragma omp target update to(dptr[0:2][1:3][0:4])
   }
 };

clang/test/OpenMP/target_update_strided_ptr_messages_from.c

@@ -0,0 +1,40 @@
+// 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;
+  double *data;
+
+  // Valid strided array sections with FROM
+  #pragma omp target update from(data[0:8:2]) // OK - even indices
+  {}
+
+  #pragma omp target update from(data[1:4:3]) // OK - odd start with stride
+  {}
+
+  #pragma omp target update from(data[2:3:5]) // OK - large stride
+  {}
+
+  // Missing stride (default = 1)
+  #pragma omp target update from(data[0:8]) // OK - default stride
+  {}
+
+  #pragma omp target update from(data[4:len-4]) // OK - computed length
+  {}
+
+  // Invalid stride expressions
+  #pragma omp target update from(data[0:8: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'}}
+
+  #pragma omp target update from(data[1:5:-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'}}
+
+  // Syntax errors
+  #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'}}
+  {}
+
+  #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_ptr_messages_to.c

@@ -0,0 +1,40 @@
+// 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;
+  double *data;
+
+  // Valid strided array sections with TO
+  #pragma omp target update to(data[0:8:2]) // OK - even indices
+  {}
+
+  #pragma omp target update to(data[1:4:3]) // OK - odd start with stride
+  {}
+
+  #pragma omp target update to(data[2:3:5]) // OK - large stride
+  {}
+
+  // Missing stride (default = 1)
+  #pragma omp target update to(data[0:8]) // OK - default stride
+  {}
+
+  #pragma omp target update to(data[4:len-4]) // OK - computed length
+  {}
+
+  // Invalid stride expressions
+  #pragma omp target update to(data[0:8: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 to(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'}}
+
+  #pragma omp target update to(data[1:5:-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'}}
+
+  // Syntax errors
+  #pragma omp target update to(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'}}
+  {}
+
+  #pragma omp target update to(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_ptr_multiple_messages_from.c

@@ -0,0 +1,52 @@
+// 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 = 12;
+  double *data;
+  double *data1;
+  double *data2;
+
+  // Valid multiple strided array sections
+  #pragma omp target update from(data1[0:6:2], data2[0:4:3]) // OK - different strides
+  {}
+
+  #pragma omp target update from(data1[1:2:3], data2[2:3:2]) // OK - with offsets
+  {}
+
+  // Mixed strided and regular sections
+  #pragma omp target update from(data1[0:len], data2[0:4:2]) // OK - mixed
+  {}
+
+  #pragma omp target update from(data1[1:3:2], data2[0:len]) // OK - reversed mix
+  {}
+
+  // Using the single data pointer with strides
+  #pragma omp target update from(data[0:4:2]) // OK - single pointer
+  {}
+
+  // 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}}
+
+  #pragma omp target update from(data1[0:3:-1], data2[0:2:2]) // expected-error {{section stride is evaluated to a non-positive value -1}}
+
+  #pragma omp target update from(data[0:4:0], data1[0:2:1]) // 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[1:4:stride1+1], data2[0:3:stride2-1]) // OK - expressions
+  {}
+
+  // Mix all three pointers
+  #pragma omp target update from(data[0:2:3], data1[1:3:2], data2[2:2:4]) // OK - three arrays
+  {}
+
+  // Syntax errors in multiple arrays
+  #pragma omp target update from(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+  #pragma omp target update from(data1[0:4:2:3], data2[0:3:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+  #pragma omp target update from(data[0:4:2], data1[0:3:2:1], data2[0:2:3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+  return 0;
+}

clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c

@@ -0,0 +1,52 @@
+// 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 = 12;
+  double *data;
+  double *data1;
+  double *data2;
+
+  // Valid multiple strided array sections
+  #pragma omp target update to(data1[0:6:2], data2[0:4:3]) // OK - different strides
+  {}
+
+  #pragma omp target update to(data1[1:2:3], data2[2:3:2]) // OK - with offsets
+  {}
+
+  // Mixed strided and regular sections
+  #pragma omp target update to(data1[0:len], data2[0:4:2]) // OK - mixed
+  {}
+
+  #pragma omp target update to(data1[1:3:2], data2[0:len]) // OK - reversed mix
+  {}
+
+  // Using the single data pointer with strides
+  #pragma omp target update to(data[0:4:2]) // OK - single pointer
+  {}
+
+  // Invalid stride in one of multiple sections
+  #pragma omp target update to(data1[0:3:4], data2[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+
+  #pragma omp target update to(data1[0:3:-1], data2[0:2:2]) // expected-error {{section stride is evaluated to a non-positive value -1}}
+
+  #pragma omp target update to(data[0:4:0], data1[0:2:1]) // 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 to(data1[1:4:stride1+1], data2[0:3:stride2-1]) // OK - expressions
+  {}
+
+  // Mix all three pointers
+  #pragma omp target update to(data[0:2:3], data1[1:3:2], data2[2:2:4]) // OK - three arrays
+  {}
+
+  // Syntax errors in multiple arrays
+  #pragma omp target update to(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+  #pragma omp target update to(data1[0:4:2:3], data2[0:3:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+  #pragma omp target update to(data[0:4:2], data1[0:3:2:1], data2[0:2:3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+  return 0;
+}

clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c

@@ -0,0 +1,37 @@
+// 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 = 11;
+  double *data;
+
+  // Valid partial strided sections with FROM
+  #pragma omp target update from(data[0:2:3]) // OK - partial coverage
+  {}
+
+  #pragma omp target update from(data[1:3:4]) // OK - offset with partial stride
+  {}
+
+  #pragma omp target update from(data[2:2:5]) // OK - large partial stride
+  {}
+
+  // Stride larger than remaining elements
+  #pragma omp target update from(data[0:2:10]) // OK - stride > array size
+  {}
+
+  #pragma omp target update from(data[0:3:len]) // OK - stride = len
+  {}
+
+  // Complex expressions
+  int offset = 1;
+  int stride = 2;
+
+  // Runtime-dependent invalid strides
+  #pragma omp target update from(data[0:4:offset-1]) // OK if offset > 1
+  {}
+
+  // Compile-time invalid strides
+  #pragma omp target update from(data[1:2:-3]) // expected-error {{section stride is evaluated to a non-positive value -3}} 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_ptr_partial_messages_to.c

@@ -0,0 +1,36 @@
+// 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 = 11;
+  double *data;
+
+  // Valid partial strided sections with TO
+  #pragma omp target update to(data[0:2:3]) // OK - partial coverage
+  {}
+
+  #pragma omp target update to(data[1:3:4]) // OK - offset with partial stride
+  {}
+
+  #pragma omp target update to(data[2:2:5]) // OK - large partial stride
+  {}
+
+  // Stride larger than remaining elements
+  #pragma omp target update to(data[0:2:10]) // OK - stride > array size
+  {}
+
+  #pragma omp target update to(data[0:3:len]) // OK - stride = len
+  {}
+
+  int offset = 1;
+  int stride = 2;
+
+  // Runtime-dependent invalid strides
+  #pragma omp target update to(data[0:4:offset-1]) // OK if offset > 1
+  {}
+
+  // Compile-time invalid strides
+  #pragma omp target update to(data[1:2:-3]) // expected-error {{section stride is evaluated to a non-positive value -3}} 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_struct_messages_from.c

@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+#define N 8
+typedef struct { 
+  double data[N]; 
+  int len; 
+} T;
+
+int main(int argc, char **argv) {
+  T s;
+  s.len = N;
+
+  // Valid strided array sections
+  #pragma omp target update from(s.data[0:4:2]) // OK
+  {}
+
+  #pragma omp target update from(s.data[1:3:2]) // OK
+  {}
+
+  // Missing stride (default = 1)
+  #pragma omp target update from(s.data[0:4]) // OK
+  {}
+
+  // Invalid stride expressions
+  #pragma omp target update from(s.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(s.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(s.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(s.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;
+}