cchen created this revision.
cchen added a reviewer: ABataev.
Herald added subscribers: cfe-commits, sstefan1, guansong, yaxunl.
Herald added a reviewer: jdoerfert.
Herald added a project: clang.
In order not to modify the `tgt_target_data_update` information but still be
able to pass the extra information for non-contiguous map item (offset,
count, and stride for each dimension), this patch overload `arg` when
the maptype is set as `OMP_MAP_DESCRIPTOR`. The origin `arg` is for
passing the pointer information, however, the overloaded `arg` is an
array of descriptor_dim:
struct descriptor_dim {
int64_t offset;
int64_t count;
int64_t stride
};
and the array size is the same as dimension size. In addition, since we
have count and stride information in descriptor_dim, we can replace/overload the
`arg_size` parameter by using dimension size.
For supporting `stride` in array section, we use a dummy dimension in
descriptor to store the unit size. The formula for counting the stride
in dimension D_n: `unit size * (D_0 * D_1 ... * D_n-1) * D_n.stride`.
Demonstrate how it works:
double arr[3][4][5];
D0: { offset = 0, count = 1, stride = 8 } //
offset, count, dimension size always be 0, 1, 1 for this extra dimension,
stride is the unit size
D1: { offset = 0, count = 2, stride = 8 * 1 * 2 = 16 } //
stride = unit size * (product of dimension size of D0) * D1.stride = 4 * 1 * 2
= 8
D2: { offset = 2, count = 2, stride = 8 * (1 * 5) * 1 = 40 } //
stride = unit size * (product of dimension size of D0, D1) * D2.stride = 4 * 5
* 1 = 20
D3: { offset = 0, count = 2, stride = 8 * (1 * 5 * 4) * 2 = 320 } //
stride = unit size * (product of dimension size of D0, D1, D2) * D3.stride = 4
* 25 * 2 = 200
// X here means we need to offload this data, therefore, runtime will transfer
// data from offset 80, 96, 120, 136, 400, 416, 440, 456
// Runtime patch: https://reviews.llvm.org/D82245
// OOOOO OOOOO OOOOO
// OOOOO OOOOO OOOOO
// XOXOO OOOOO XOXOO
// XOXOO OOOOO XOXOO
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D84192
Files:
clang/include/clang/AST/OpenMPClause.h
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntime.h
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/OpenMP/target_update_ast_print.cpp
clang/test/OpenMP/target_update_codegen.cpp
clang/test/OpenMP/target_update_messages.cpp
clang/test/OpenMP/target_update_to_messages.cpp
Index: clang/test/OpenMP/target_update_to_messages.cpp
===================================================================
--- clang/test/OpenMP/target_update_to_messages.cpp
+++ clang/test/OpenMP/target_update_to_messages.cpp
@@ -79,6 +79,10 @@
#pragma omp target update to(*(*(this->ptr)+a+this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
#pragma omp target update to(*(this+this)) // expected-error {{invalid operands to binary expression ('S8 *' and 'S8 *')}}
{}
+
+ double marr[10][5][10];
+#pragma omp target update to(marr [0:1][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
}
};
Index: clang/test/OpenMP/target_update_messages.cpp
===================================================================
--- clang/test/OpenMP/target_update_messages.cpp
+++ clang/test/OpenMP/target_update_messages.cpp
@@ -1,8 +1,8 @@
-// RUN: %clang_cc1 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp50 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wuninitialized
void xxx(int argc) {
int x; // expected-note {{initialize the variable 'x' to silence this warning}}
@@ -39,10 +39,33 @@
foo();
}
+ double marr[10][5][10];
+#pragma omp target update to(marr[0:2][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+#pragma omp target update from(marr[0:2][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+#pragma omp target update to(marr[0:][1:2:2][1:2]) // le50-error {{array section does not specify length for outermost dimension}} le45-error {{expected ']'}} le45-note {{to match this '['}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+#pragma omp target update from(marr[0:][1:2:2][1:2]) // le50-error {{array section does not specify length for outermost dimension}} le45-error {{expected ']'}} le45-note {{to match this '['}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ int arr[4][3][2][1];
+#pragma omp target update to(arr[0:2][2:4][:2][1]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+#pragma omp target update from(arr[0:2][2:4][:2][1]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ double ***dptr;
+#pragma omp target update to(dptr[0:2][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+#pragma omp target update from(dptr[0:2][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
int iarr[5][5];
-#pragma omp target update to(iarr[0:][1:2:-1]) // omp50-error {{section stride is evaluated to a non-positive value -1}} omp45-error {{expected ']'}} omp45-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(iarr[0:][1:2:-1]) // le50-error {{section stride is evaluated to a non-positive value -1}} le45-error {{expected ']'}} le45-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(iarr[0:][1:2:-1]) // le50-error {{section stride is evaluated to a non-positive value -1}} le45-error {{expected ']'}} le45-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(iarr[0: :2][1:2]) // le50-error {{array section does not specify length for outermost dimension}} le45-error {{expected ']'}} le45-note {{to match this '['}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le45-error {{expected expression}}
{}
-#pragma omp target update from(iarr[0:][1:2:-1]) // omp50-error {{section stride is evaluated to a non-positive value -1}} omp45-error {{expected ']'}} omp45-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(iarr[0: :2][1:2]) // le50-error {{array section does not specify length for outermost dimension}} le45-error {{expected ']'}} le45-note {{to match this '['}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le45-error {{expected expression}}
return tmain(argc, argv);
}
Index: clang/test/OpenMP/target_update_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_update_codegen.cpp
+++ clang/test/OpenMP/target_update_codegen.cpp
@@ -1059,5 +1059,495 @@
#pragma omp target update from(([sa][5])f)
}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64
+// RUN: %clang_cc1 -DCK19 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-32
+// RUN: %clang_cc1 -DCK19 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-32
+
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK19 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK19 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK19
+
+// CK19: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
+
+// CK19: [[MSIZE:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4, i64 3]
+// CK19: [[MTYPE:@.+]] = {{.+}}constant [3 x i64] [i64 17592186044449, i64 33, i64 17592186044449]
+
+// CK19-LABEL: _Z3foo
+void foo(int arg) {
+ int arr[3][4][5], x;
+ float farr[4][3];
+
+ // CK19: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+ // CK19: [[DIMS_2:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]],
+ // CK19: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [5 x i32]]], [3 x [4 x [5 x i32]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+ // CK19: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [5 x i32]], [4 x [5 x i32]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+ // CK19: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_DECAY]], {{.+}}
+ // CK19: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+ // CK19: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 1
+ // CK19: [[SUB:%.+]] = sub nuw i64 4, [[ARG_ADDR:%.+]]
+ // CK19: [[LEN:%.+]] = udiv {{.+}} [[SUB]], 1
+ // CK19: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds [4 x [3 x float]], [4 x [3 x float]]* [[FARR:%.+]], {{.+}} 0, {{.+}} 0
+ // CK19: [[ARRAY_DECAY_5:%.+]] = getelementptr inbounds [3 x float], [3 x float]* [[ARRAY_IDX_4]], {{.+}} 0, {{.+}} 0
+ // CK19: [[ARRAY_IDX_6:%.+]] = getelementptr inbounds float, float* [[ARRAY_DECAY_5:%.+]], {{.+}} 1
+ // CK19: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i{{.+}} 0, i{{.+}} 0
+ // CK19: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+ // CK19: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+ // CK19: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 0, i64* [[OFFSET]],
+ // CK19: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 2, i64* [[COUNT]],
+ // CK19: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 80, i64* [[STRIDE]],
+ // CK19: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+ // CK19: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 [[ARG:%.+]], i64* [[OFFSET_2]],
+ // CK19: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 [[LEN]], i64* [[COUNT_2]],
+ // CK19: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 20, i64* [[STRIDE_2]],
+ // CK19: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+ // CK19: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 1, i64* [[OFFSET_3]],
+ // CK19: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 4, i64* [[COUNT_3]],
+ // CK19: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 4, i64* [[STRIDE_3]],
+ // CK19: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+ // CK19: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 0, i64* [[OFFSET_4]],
+ // CK19: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 1, i64* [[COUNT_4]],
+ // CK19: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 4, i64* [[STRIDE_4]],
+ // CK19: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+ // CK19: [[PTRS:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0
+ // CK19: store i8* [[PC0]], i8** [[PTRS]],
+ // CK19: [[DIM_5:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 0
+ // CK19: [[OFFSET_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 0, i64* [[OFFSET_2_1]],
+ // CK19: [[COUNT_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 2, i64* [[COUNT_2_1]],
+ // CK19: [[STRIDE_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 12, i64* [[STRIDE_2_1]],
+ // CK19: [[DIM_6:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 1
+ // CK19: [[OFFSET_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 1, i64* [[OFFSET_2_2]],
+ // CK19: [[COUNT_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 2, i64* [[COUNT_2_2]],
+ // CK19: [[STRIDE_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 4, i64* [[STRIDE_2_2]],
+ // CK19: [[DIM_7:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 2
+ // CK19: [[OFFSET_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 0, i64* [[OFFSET_2_3]],
+ // CK19: [[COUNT_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 1, i64* [[COUNT_2_3]],
+ // CK19: [[STRIDE_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 4, i64* [[STRIDE_2_3]],
+ // CK19: [[PC1:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]] to i8*
+ // CK19: [[PTRS_2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 2
+ // CK19: store i8* [[PC1]], i8** [[PTRS_2]],
+ // CK19-DAG: call void @__tgt_target_data_update(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE]]{{.+}})
+ // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+#pragma omp target update to(arr[0:2][arg:][1:4], x, farr[0:2][1:2])
+ { ++arg; }
+}
+
+#endif
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-64
+// RUN: %clang_cc1 -DCK20 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-64
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-32
+// RUN: %clang_cc1 -DCK20 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-32
+
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK20
+
+struct ST {
+ int a;
+ double *b;
+};
+
+// CK20: [[STRUCT_ST:%.+]] = type { i32, double* }
+// CK20: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
+
+// CK20: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 3]
+// CK20: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK20-LABEL: _Z3foo
+void foo(int arg) {
+ ST arr[3][4];
+ // CK20: [[DIMS:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]],
+ // CK20: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [[STRUCT_ST]]]], [3 x [4 x [[STRUCT_ST]]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+ // CK20: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [[STRUCT_ST]]], [4 x [[STRUCT_ST]]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+ // CK20: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [[STRUCT_ST]], [[STRUCT_ST]]* [[ARRAY_DECAY]], {{.+}}
+ // CK20: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+ // CK20: [[BPC:%.+]] = bitcast i8** [[BP0]] to [3 x [4 x [[STRUCT_ST]]]]**
+ // CK20: store [3 x [4 x [[STRUCT_ST]]]]* [[ARR]], [3 x [4 x [[STRUCT_ST]]]]** [[BPC]],
+ // CK20: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0
+ // CK20: [[PC:%.+]] = bitcast i8** [[P0]] to [[STRUCT_ST]]**
+ // CK20: store [[STRUCT_ST]]* [[ARRAY_IDX_1]], [[STRUCT_ST]]** [[PC]],
+ // CK20: [[DIM_1:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+ // CK20: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+ // CK20: store i64 0, i64* [[OFFSET]],
+ // CK20: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+ // CK20: store i64 2, i64* [[COUNT]],
+ // CK20: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+ // CK20: store i64 {{32|64}}, i64* [[STRIDE]],
+ // CK20: [[DIM_2:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+ // CK20: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+ // CK20: store i64 1, i64* [[OFFSET_2]],
+ // CK20: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+ // CK20: store i64 4, i64* [[COUNT_2]],
+ // CK20: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+ // CK20: store i64 {{8|16}}, i64* [[STRIDE_2]],
+ // CK20: [[DIM_3:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+ // CK20: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+ // CK20: store i64 0, i64* [[OFFSET_3]],
+ // CK20: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+ // CK20: store i64 1, i64* [[COUNT_3]],
+ // CK20: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+ // CK20: store i64 {{8|16}}, i64* [[STRIDE_3]],
+ // CK20-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+ // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK20-DAG: [[PC0:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+ // CK20-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+ // CK20-DAG: store i8* [[PC0]], i8** [[PTRS]],
+
+#pragma omp target update to(arr[0:2][1:4])
+ { ++arg; }
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-64
+// RUN: %clang_cc1 -DCK21 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-64
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-32
+// RUN: %clang_cc1 -DCK21 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-32
+
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK21
+
+// CK21: [[STRUCT_ST:%.+]] = type { [10 x [10 x [10 x double*]]] }
+// CK21: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
+
+// CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 299067162755073]
+
+struct ST {
+ double *dptr[10][10][10];
+
+ // CK21: _ZN2ST3fooEv
+ void foo() {
+ // CK21: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+ // CK21: [[ARRAY_IDX:%.+]] = getelementptr inbounds [10 x [10 x [10 x double*]]], [10 x [10 x [10 x double*]]]* [[DPTR:%.+]], {{.+}} 0, {{.+}} 0
+ // CK21: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [10 x [10 x double*]], [10 x [10 x double*]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+ // CK21: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [10 x double*], [10 x double*]* [[ARRAY_DECAY]], {{.+}} 1
+ // CK21: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [10 x double*], [10 x double*]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+ // CK21: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 0
+ // CK21: [[BP0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+ // CK21: [[P0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+ // CK21: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+ // CK21: [[OFFSET_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+ // CK21: store i64 0, i64* [[OFFSET_1]],
+ // CK21: [[COUNT_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+ // CK21: store i64 2, i64* [[COUNT_1]],
+ // CK21: [[STRIDE_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+ // CK21: store i64 {{400|800}}, i64* [[STRIDE_1]],
+ // CK21: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+ // CK21: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+ // CK21: store i64 1, i64* [[OFFSET_2]],
+ // CK21: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+ // CK21: store i64 3, i64* [[COUNT_2]],
+ // CK21: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+ // CK21: store i64 {{40|80}}, i64* [[STRIDE_2]],
+ // CK21: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+ // CK21: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+ // CK21: store i64 0, i64* [[OFFSET_3]],
+ // CK21: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+ // CK21: store i64 4, i64* [[COUNT_3]],
+ // CK21: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+ // CK21: store i64 {{4|8}}, i64* [[STRIDE_3]],
+ // CK21: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+ // CK21: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+ // CK21: store i64 0, i64* [[OFFSET_4]],
+ // CK21: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+ // CK21: store i64 1, i64* [[COUNT_4]],
+ // CK21: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+ // CK21: store i64 {{4|8}}, i64* [[STRIDE_4]],
+ // CK21-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}}* [[GEPSZ:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE]]{{.+}})
+ // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK21-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+ // CK21-DAG: [[PTRS:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i32 0, i32 0
+ // CK21-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(dptr[0:2][1:3][0:4])
+ }
+};
+
+void bar() {
+ ST st;
+ st.foo();
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-64
+// RUN: %clang_cc1 -DCK22 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-64
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-32
+// RUN: %clang_cc1 -DCK22 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-32
+
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK22
+
+// CK22: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
+
+// CK22: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK22: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+struct ST {
+ // CK22: _ZN2ST3fooEPA10_Pi
+ void foo(int *arr[5][10]) {
+ // CK22: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+ // CK22: [[ARRAY_IDX:%.+]] = getelementptr inbounds [10 x i32*], [10 x i32*]* [[ARR:%.+]], {{.+}} 0
+ // CK22: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [10 x i32*], [10 x i32*]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+ // CK22: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds i32*, i32** [[ARRAY_DECAY:%.+]], {{.+}} 1
+ // CK22: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+ // CK22: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+ // CK22: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+ // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+ // CK22: store i64 0, i64* [[OFFSET]],
+ // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+ // CK22: store i64 2, i64* [[COUNT]],
+ // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+ // CK22: store i64 200, i64* [[STRIDE]],
+ // CK22: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+ // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+ // CK22: store i64 1, i64* [[OFFSET]],
+ // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+ // CK22: store i64 3, i64* [[COUNT]],
+ // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+ // CK22: store i64 40, i64* [[STRIDE]],
+ // CK22: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+ // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+ // CK22: store i64 0, i64* [[OFFSET]],
+ // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+ // CK22: store i64 4, i64* [[COUNT]],
+ // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+ // CK22: store i64 4, i64* [[STRIDE]],
+ // CK22: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+ // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+ // CK22: store i64 0, i64* [[OFFSET]],
+ // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+ // CK22: store i64 1, i64* [[COUNT]],
+ // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+ // CK22: store i64 4, i64* [[STRIDE]],
+ // CK22-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+ // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK22-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+ // CK22-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+ // CK22-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(arr[0:2][1:3][0:4])
+ }
+};
+
+void bar() {
+ ST st;
+ int *arr[5][10];
+ st.foo(arr);
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK23 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64
+// RUN: %clang_cc1 -DCK23 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64
+// RUN: %clang_cc1 -DCK23 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-32
+// RUN: %clang_cc1 -DCK23 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-32
+
+// RUN: %clang_cc1 -DCK23 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK23 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK23 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK23 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK23
+
+// CK23: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
+
+// CK23: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK23: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK23: foo
+void foo(int arg) {
+ float farr[5][5][5];
+ // CK23: [[ARG_ADDR:%.+]] = alloca i32,
+ // CK23: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+ // CK23: [[ARRAY_IDX:%.+]] = getelementptr inbounds [5 x [5 x [5 x float]]], [5 x [5 x [5 x float]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+ // CK23: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [5 x [5 x float]], [5 x [5 x float]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+ // CK23: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x float], [5 x float]* [[ARRAY_DECAY]], {{.+}}
+ // CK23: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x float], [5 x float]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+ // CK23: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds float, float* [[ARRAY_DECAY_2]], {{.+}}
+ // CK23: [[MUL:%.+]] = mul nuw i64 4,
+ // CK23: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+ // CK23: [[BPC:%.+]] = bitcast i8** [[BP0]] to [5 x [5 x [5 x float]]]**
+ // CK23: store [5 x [5 x [5 x float]]]* [[ARR]], [5 x [5 x [5 x float]]]** [[BPC]],
+ // CK23: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0
+ // CK23: [[PC:%.+]] = bitcast i8** [[P0]] to float**
+ // CK23: store float* [[ARRAY_IDX_2]], float** [[PC]],
+ // CK23: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+ // CK23: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+ // CK23: store i64 0, i64* [[OFFSET]],
+ // CK23: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+ // CK23: store i64 2, i64* [[COUNT]],
+ // CK23: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+ // CK23: store i64 200, i64* [[STRIDE]],
+ // CK23: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+ // CK23: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+ // CK23: store i64 1, i64* [[OFFSET_2]],
+ // CK23: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+ // CK23: store i64 2, i64* [[COUNT_2]],
+ // CK23: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+ // CK23: store i64 20, i64* [[STRIDE_2]],
+ // CK23: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+ // CK23: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+ // CK23: store i64 0, i64* [[OFFSET_3]],
+ // CK23: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+ // CK23: store i64 2, i64* [[COUNT_3]],
+ // CK23: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+ // CK23: store i64 [[MUL]], i64* [[STRIDE_3]],
+ // CK23: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+ // CK23: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+ // CK23: store i64 0, i64* [[OFFSET_4]],
+ // CK23: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+ // CK23: store i64 1, i64* [[COUNT_4]],
+ // CK23: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+ // CK23: store i64 4, i64* [[STRIDE_4]],
+ // CK23-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+ // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK23-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+ // CK23-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+ // CK23-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(farr[0:2:2][1:2:1][0:2:arg])
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-64
+// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-64
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-32
+// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-32
+
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK24 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK24 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK24
+
+// CK24: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
+
+// CK24: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK24: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK24: foo
+void foo(int arg) {
+ double darr[3][4][5];
+ // CK24: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+ // CK24: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [5 x double]]], [3 x [4 x [5 x double]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+ // CK24: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [5 x double]], [4 x [5 x double]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+ // CK24: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x double], [5 x double]* [[ARRAY_DECAY]], {{.+}}
+ // CK24: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x double], [5 x double]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+ // CK24: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds double, double* [[ARRAY_DECAY_2]], {{.+}}
+ // CK24: [[MUL:%.+]] = mul nuw i64 8,
+ // CK24: [[SUB:%.+]] = sub nuw i64 4, [[ARG:%.+]]
+ // CK24: [[LEN:%.+]] = udiv {{.+}} [[SUB]], 1
+ // CK24: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+ // CK24: [[BPC:%.+]] = bitcast i8** [[BP0]] to [3 x [4 x [5 x double]]]**
+ // CK24: store [3 x [4 x [5 x double]]]* [[ARR]], [3 x [4 x [5 x double]]]** [[BPC]],
+ // CK24: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0
+ // CK24: [[PC:%.+]] = bitcast i8** [[P0]] to double**
+ // CK24: store double* [[ARRAY_IDX_2]], double** [[PC]],
+ // CK24: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+ // CK24: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+ // CK24: store i64 0, i64* [[OFFSET]],
+ // CK24: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+ // CK24: store i64 2, i64* [[COUNT]],
+ // CK24: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+ // CK24: store i64 320, i64* [[STRIDE]],
+ // CK24: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+ // CK24: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+ // CK24: store i64 [[ARG]], i64* [[OFFSET_2]],
+ // CK24: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+ // CK24: store i64 [[LEN]], i64* [[COUNT_2]],
+ // CK24: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+ // CK24: store i64 40, i64* [[STRIDE_2]],
+ // CK24: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+ // CK24: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+ // CK24: store i64 0, i64* [[OFFSET_3]],
+ // CK24: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+ // CK24: store i64 2, i64* [[COUNT_3]],
+ // CK24: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+ // CK24: store i64 [[MUL]], i64* [[STRIDE_3]],
+ // CK24: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+ // CK24: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+ // CK24: store i64 0, i64* [[OFFSET_4]],
+ // CK24: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+ // CK24: store i64 1, i64* [[COUNT_4]],
+ // CK24: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+ // CK24: store i64 8, i64* [[STRIDE_4]],
+ // CK24-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+ // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK24-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+ // CK24-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+ // CK24-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(darr[0:2:2][arg: :1][:2:arg])
+}
+
#endif
#endif
Index: clang/test/OpenMP/target_update_ast_print.cpp
===================================================================
--- clang/test/OpenMP/target_update_ast_print.cpp
+++ clang/test/OpenMP/target_update_ast_print.cpp
@@ -5,6 +5,14 @@
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -DOMP5 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5
+// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5
+
+// RUN: %clang_cc1 -DOMP5 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5
+// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5
// expected-no-diagnostics
#ifndef HEADER
@@ -21,10 +29,61 @@
#pragma omp target update from(b, ([a][targ])p) if(l<5) device(l-1) nowait depend(inout:l)
+#ifdef OMP5
+ U marr[10][10][10];
+#pragma omp target update to(marr[2][0:2][0:2])
+
+#pragma omp target update from(marr[2][0:2][0:2])
+
+#pragma omp target update from(marr[:2][0:2][0:2:1])
+
+#pragma omp target update to(marr[:l][:l][l:])
+
+#pragma omp target update to(marr[:2][:1][:])
+
+#pragma omp target update from(marr[:2][:1][:])
+
+#pragma omp target update to(marr[:2][:][:1])
+
+#pragma omp target update from(marr[:2][:][:1])
+
+#pragma omp target update to(marr[:2][:] [1:])
+
+#pragma omp target update from(marr[:2][:][1:])
+
+#pragma omp target update to(marr[:1][3:2][:2])
+
+#pragma omp target update from(marr[:1][3:2][:2])
+
+#pragma omp target update to(marr[:1][:2][0])
+
+#pragma omp target update from(marr[:1][:2][0])
+
int arr[100][100];
+
#pragma omp target update to(arr[2][0:1:2])
#pragma omp target update from(arr[2][0:1:2])
+
+// OMP5: marr[10][10][10];
+// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][0:2][0:2:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:l][:l][l:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0])
+// OMP5-NEXT: int arr[100][100];
+// OMP5-NEXT: #pragma omp target update to(arr[2][0:1:2])
+// OMP5-NEXT: #pragma omp target update from(arr[2][0:1:2])
+#endif
return a + targ + (T)b;
}
// CHECK: static T a, *p;
@@ -42,9 +101,6 @@
// CHECK-NEXT: int l;
// CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l)
// CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
-// CHECK: int arr[100][100];
-// CHECK-NEXT: #pragma omp target update to(arr[2][0:1:2])
-// CHECK-NEXT: #pragma omp target update from(arr[2][0:1:2])
int main(int argc, char **argv) {
static int a;
@@ -58,10 +114,47 @@
// CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n) nowait depend(in : n)
#pragma omp target update from(f) if(f<0.0) device(n+1) nowait depend(in:n)
// CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1) nowait depend(in : n)
+
+#ifdef OMP5
+float marr[10][10][10];
+// OMP5: marr[10][10][10];
+#pragma omp target update to(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
+#pragma omp target update from(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
+#pragma omp target update to(marr[:n][:n][n:])
+// OMP5: #pragma omp target update to(marr[:n][:n][n:])
+#pragma omp target update from(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:])
+#pragma omp target update to(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1])
+#pragma omp target update from(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1])
+#pragma omp target update to(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:])
+#pragma omp target update from(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:])
+#pragma omp target update to(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
+#pragma omp target update from(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
+#pragma omp target update to(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0])
+#pragma omp target update from(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0])
+#pragma omp target update to(marr[:2:][0:2][0:2:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:2:][0:2][0:2:1])
+#pragma omp target update from(marr[:2:][0:2][0:2:1])
+// OMP5-NEXT: #pragma omp target update from(marr[:2:][0:2][0:2:1])
+#pragma omp target update to(marr[:2:][:2:][0:2:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:2:][:2:][0:2:1])
+#pragma omp target update from(marr[:2:][:2:][0:2:1])
+// OMP5-NEXT: #pragma omp target update from(marr[:2:][:2:][0:2:1])
#pragma omp target update to(argv[2][0:1:2])
-// CHECK-NEXT: #pragma omp target update to(argv[2][0:1:2])
+// OMP5-NEXT: #pragma omp target update to(argv[2][0:1:2])
#pragma omp target update from(argv[2][0:1:2])
-// CHECK-NEXT: #pragma omp target update from(argv[2][0:1:2])
+// OMP5-NEXT: #pragma omp target update from(argv[2][0:1:2])
+#endif
return foo(argc, f) + foo(argv[0][0], f) + a;
}
Index: clang/lib/Serialization/ASTWriter.cpp
===================================================================
--- clang/lib/Serialization/ASTWriter.cpp
+++ clang/lib/Serialization/ASTWriter.cpp
@@ -6597,6 +6597,7 @@
Record.push_back(N);
for (auto &M : C->all_components()) {
Record.AddStmt(M.getAssociatedExpression());
+ Record.writeBool(M.isNonContiguous());
Record.AddDeclRef(M.getAssociatedDeclaration());
}
}
@@ -6621,6 +6622,7 @@
Record.push_back(N);
for (auto &M : C->all_components()) {
Record.AddStmt(M.getAssociatedExpression());
+ Record.writeBool(M.isNonContiguous());
Record.AddDeclRef(M.getAssociatedDeclaration());
}
}
Index: clang/lib/Serialization/ASTReader.cpp
===================================================================
--- clang/lib/Serialization/ASTReader.cpp
+++ clang/lib/Serialization/ASTReader.cpp
@@ -12531,10 +12531,10 @@
SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
Components.reserve(TotalComponents);
for (unsigned i = 0; i < TotalComponents; ++i) {
- Expr *AssociatedExpr = Record.readExpr();
+ Expr *AssociatedExprPr = Record.readExpr();
auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
- Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
- AssociatedExpr, AssociatedDecl));
+ Components.emplace_back(AssociatedExprPr, AssociatedDecl,
+ /*IsNonContiguous=*/false);
}
C->setComponents(Components, ListSizes);
}
@@ -12648,10 +12648,10 @@
SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
Components.reserve(TotalComponents);
for (unsigned i = 0; i < TotalComponents; ++i) {
- Expr *AssociatedExpr = Record.readSubExpr();
+ Expr *AssociatedExprPr = Record.readSubExpr();
+ bool IsNonContiguous = Record.readBool();
auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
- Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
- AssociatedExpr, AssociatedDecl));
+ Components.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous);
}
C->setComponents(Components, ListSizes);
}
@@ -12698,10 +12698,10 @@
SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
Components.reserve(TotalComponents);
for (unsigned i = 0; i < TotalComponents; ++i) {
- Expr *AssociatedExpr = Record.readSubExpr();
+ Expr *AssociatedExprPr = Record.readSubExpr();
+ bool IsNonContiguous = Record.readBool();
auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
- Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
- AssociatedExpr, AssociatedDecl));
+ Components.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous);
}
C->setComponents(Components, ListSizes);
}
@@ -12748,10 +12748,10 @@
SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
Components.reserve(TotalComponents);
for (unsigned i = 0; i < TotalComponents; ++i) {
- Expr *AssociatedExpr = Record.readSubExpr();
+ auto *AssociatedExprPr = Record.readSubExpr();
auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
- Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
- AssociatedExpr, AssociatedDecl));
+ Components.emplace_back(AssociatedExprPr, AssociatedDecl,
+ /*IsNonContiguous=*/false);
}
C->setComponents(Components, ListSizes);
}
@@ -12792,8 +12792,8 @@
for (unsigned i = 0; i < TotalComponents; ++i) {
Expr *AssociatedExpr = Record.readSubExpr();
auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
- Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
- AssociatedExpr, AssociatedDecl));
+ Components.emplace_back(AssociatedExpr, AssociatedDecl,
+ /*IsNonContiguous*/ false);
}
C->setComponents(Components, ListSizes);
}
@@ -12835,8 +12835,8 @@
for (unsigned i = 0; i < TotalComponents; ++i) {
Expr *AssociatedExpr = Record.readSubExpr();
auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
- Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
- AssociatedExpr, AssociatedDecl));
+ Components.emplace_back(AssociatedExpr, AssociatedDecl,
+ /*IsNonContiguous=*/false);
}
C->setComponents(Components, ListSizes);
}
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -48,7 +48,7 @@
static const Expr *checkMapClauseExpressionBase(
Sema &SemaRef, Expr *E,
OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
- OpenMPClauseKind CKind, bool NoDiagnose);
+ OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose);
namespace {
/// Default data sharing attributes, which can be applied to directive.
@@ -3535,6 +3535,7 @@
if (isOpenMPTargetExecutionDirective(DKind)) {
OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map,
+ Stack->getCurrentDirective(),
/*NoDiagnose=*/true))
return;
const auto *VD = cast<ValueDecl>(
@@ -16597,11 +16598,14 @@
class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
Sema &SemaRef;
OpenMPClauseKind CKind = OMPC_unknown;
+ OpenMPDirectiveKind DKind = OMPD_unknown;
OMPClauseMappableExprCommon::MappableExprComponentList &Components;
+ bool IsNonContiguous = false;
bool NoDiagnose = false;
const Expr *RelevantExpr = nullptr;
bool AllowUnitySizeArraySection = true;
bool AllowWholeSizeArraySection = true;
+ bool AllowAnotherPtr = true;
SourceLocation ELoc;
SourceRange ERange;
@@ -16626,7 +16630,7 @@
assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
RelevantExpr = DRE;
// Record the component.
- Components.emplace_back(DRE, DRE->getDecl());
+ Components.emplace_back(DRE, DRE->getDecl(), IsNonContiguous);
return true;
}
@@ -16698,7 +16702,7 @@
AllowWholeSizeArraySection = false;
// Record the component.
- Components.emplace_back(ME, FD);
+ Components.emplace_back(ME, FD, IsNonContiguous);
return RelevantExpr || Visit(E);
}
@@ -16736,7 +16740,7 @@
}
// Record the component - we don't have any declaration associated.
- Components.emplace_back(AE, nullptr);
+ Components.emplace_back(AE, nullptr, IsNonContiguous);
return RelevantExpr || Visit(E);
}
@@ -16775,6 +16779,13 @@
// pointer. Otherwise, only unitary sections are accepted.
if (NotWhole || IsPointer)
AllowWholeSizeArraySection = false;
+ } else if (DKind == OMPD_target_update &&
+ SemaRef.getLangOpts().OpenMP >= 50) {
+ if (IsPointer && !AllowAnotherPtr)
+ SemaRef.Diag(ELoc, diag::err_omp_section_length_undefined)
+ << /*array of unknown bound */ 1;
+ else
+ IsNonContiguous = true;
} else if (AllowUnitySizeArraySection && NotUnity) {
// A unity or whole array section is not allowed and that is not
// compatible with the properties of the current array section.
@@ -16784,6 +16795,9 @@
return false;
}
+ if (IsPointer)
+ AllowAnotherPtr = false;
+
if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
Expr::EvalResult ResultR;
Expr::EvalResult ResultL;
@@ -16809,14 +16823,14 @@
}
// Record the component - we don't have any declaration associated.
- Components.emplace_back(OASE, nullptr);
+ Components.emplace_back(OASE, nullptr, /*IsNonContiguous=*/false);
return RelevantExpr || Visit(E);
}
bool VisitOMPArrayShapingExpr(OMPArrayShapingExpr *E) {
Expr *Base = E->getBase();
// Record the component - we don't have any declaration associated.
- Components.emplace_back(E, nullptr);
+ Components.emplace_back(E, nullptr, IsNonContiguous);
return Visit(Base->IgnoreParenImpCasts());
}
@@ -16829,7 +16843,7 @@
}
if (!RelevantExpr) {
// Record the component if haven't found base decl.
- Components.emplace_back(UO, nullptr);
+ Components.emplace_back(UO, nullptr, /*IsNonContiguous=*/false);
}
return RelevantExpr || Visit(UO->getSubExpr()->IgnoreParenImpCasts());
}
@@ -16845,7 +16859,7 @@
// know the other subtree is just an offset)
Expr *LE = BO->getLHS()->IgnoreParenImpCasts();
Expr *RE = BO->getRHS()->IgnoreParenImpCasts();
- Components.emplace_back(BO, nullptr);
+ Components.emplace_back(BO, nullptr, false);
assert((LE->getType().getTypePtr() == BO->getType().getTypePtr() ||
RE->getType().getTypePtr() == BO->getType().getTypePtr()) &&
"Either LHS or RHS have base decl inside");
@@ -16856,7 +16870,7 @@
bool VisitCXXThisExpr(CXXThisExpr *CTE) {
assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
RelevantExpr = CTE;
- Components.emplace_back(CTE, nullptr);
+ Components.emplace_back(CTE, nullptr, IsNonContiguous);
return true;
}
bool VisitStmt(Stmt *) {
@@ -16867,10 +16881,10 @@
return RelevantExpr;
}
explicit MapBaseChecker(
- Sema &SemaRef, OpenMPClauseKind CKind,
+ Sema &SemaRef, OpenMPClauseKind CKind, OpenMPDirectiveKind DKind,
OMPClauseMappableExprCommon::MappableExprComponentList &Components,
bool NoDiagnose, SourceLocation &ELoc, SourceRange &ERange)
- : SemaRef(SemaRef), CKind(CKind), Components(Components),
+ : SemaRef(SemaRef), CKind(CKind), DKind(DKind), Components(Components),
NoDiagnose(NoDiagnose), ELoc(ELoc), ERange(ERange) {}
};
} // namespace
@@ -16882,13 +16896,31 @@
static const Expr *checkMapClauseExpressionBase(
Sema &SemaRef, Expr *E,
OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
- OpenMPClauseKind CKind, bool NoDiagnose) {
+ OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose) {
SourceLocation ELoc = E->getExprLoc();
SourceRange ERange = E->getSourceRange();
- MapBaseChecker Checker(SemaRef, CKind, CurComponents, NoDiagnose, ELoc,
+ MapBaseChecker Checker(SemaRef, CKind, DKind, CurComponents, NoDiagnose, ELoc,
ERange);
- if (Checker.Visit(E->IgnoreParens()))
+ if (Checker.Visit(E->IgnoreParens())) {
+ // Check if the highest dimension array section has length specified
+ if (SemaRef.getLangOpts().OpenMP >= 50 && !CurComponents.empty() &&
+ (CKind == OMPC_to || CKind == OMPC_from)) {
+ auto CI = CurComponents.rbegin();
+ auto CE = CurComponents.rend();
+ for (; CI != CE; ++CI) {
+ const auto *OASE =
+ dyn_cast<OMPArraySectionExpr>(CI->getAssociatedExpression());
+ if (OASE) {
+ if (!OASE->getLength())
+ SemaRef.Diag(ELoc, diag::err_array_section_does_not_specify_length)
+ << ERange;
+ else
+ break;
+ }
+ }
+ }
return Checker.getFoundBase();
+ }
return nullptr;
}
@@ -17365,7 +17397,8 @@
// Obtain the array or member expression bases if required. Also, fill the
// components array with all the components identified in the process.
const Expr *BE = checkMapClauseExpressionBase(
- SemaRef, SimpleExpr, CurComponents, CKind, /*NoDiagnose=*/false);
+ SemaRef, SimpleExpr, CurComponents, CKind, DSAS->getCurrentDirective(),
+ /*NoDiagnose=*/false);
if (!BE)
continue;
@@ -18613,8 +18646,8 @@
// only need a component.
MVLI.VarBaseDeclarations.push_back(D);
MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
- MVLI.VarComponents.back().push_back(
- OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D));
+ MVLI.VarComponents.back().emplace_back(SimpleRefExpr, D,
+ /*IsNonContiguous=*/false);
}
if (MVLI.ProcessedVarList.empty())
@@ -18665,8 +18698,8 @@
if (VD && (isa<OMPArraySectionExpr>(RefExpr->IgnoreParenImpCasts()) ||
isa<ArraySubscriptExpr>(RefExpr->IgnoreParenImpCasts())))
Component = DefaultFunctionArrayLvalueConversion(SimpleRefExpr).get();
- MVLI.VarComponents.back().push_back(
- OMPClauseMappableExprCommon::MappableComponent(Component, D));
+ MVLI.VarComponents.back().emplace_back(Component, D,
+ /*IsNonContiguous=*/false);
}
if (MVLI.ProcessedVarList.empty())
@@ -18732,7 +18765,8 @@
// Store the components in the stack so that they can be used to check
// against other clauses later on.
- OMPClauseMappableExprCommon::MappableComponent MC(SimpleRefExpr, D);
+ OMPClauseMappableExprCommon::MappableComponent MC(
+ SimpleRefExpr, D, /*IsNonContiguous=*/false);
DSAStack->addMappableExpressionComponents(
D, MC, /*WhereFoundClauseKind=*/OMPC_is_device_ptr);
Index: clang/lib/CodeGen/CGOpenMPRuntime.h
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.h
+++ clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -1625,6 +1625,14 @@
/// Map between the a declaration of a capture and the corresponding base
/// pointer address where the runtime returns the device pointers.
llvm::DenseMap<const ValueDecl *, Address> CaptureDeviceAddrMap;
+ /// The array of dimension size passed to the runtime library.
+ SmallVector<uint64_t, 4> Dims;
+ /// The array of array of offsets passed to the runtime library.
+ SmallVector<SmallVector<llvm::Value *, 4>, 4> Offsets;
+ /// The array of array of counts passed to the runtime library.
+ SmallVector<SmallVector<llvm::Value *, 4>, 4> Counts;
+ /// The array of array of strides passed to the runtime library.
+ SmallVector<SmallVector<llvm::Value *, 4>, 4> Strides;
explicit TargetDataInfo() {}
explicit TargetDataInfo(bool RequiresDevicePointerInfo)
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7040,6 +7040,10 @@
/// Close is a hint to the runtime to allocate memory close to
/// the target device.
OMP_MAP_CLOSE = 0x400,
+ /// Signal that the runtime library should use args as an array of
+ /// descriptor_dim pointers and use args_size as dims. Used when we have
+ /// non-contiguous list items in target update directive
+ OMP_MAP_DESCRIPTOR = 0x100000000000,
/// The 16 MSBs of the flags indicate whether the entry is member of some
/// struct/class.
OMP_MAP_MEMBER_OF = 0xffff000000000000,
@@ -7075,6 +7079,8 @@
using MapBaseValuesArrayTy = SmallVector<BasePointerInfo, 4>;
using MapValuesArrayTy = SmallVector<llvm::Value *, 4>;
using MapFlagsArrayTy = SmallVector<OpenMPOffloadMappingFlags, 4>;
+ using MapDimArrayTy = SmallVector<uint64_t, 4>;
+ using MapNonContiguousArrayTy = SmallVector<MapValuesArrayTy, 4>;
/// Map between a struct and the its lowest & highest elements which have been
/// mapped.
@@ -7088,6 +7094,14 @@
Address Base = Address::invalid();
};
+ struct StructNonContiguousInfo {
+ bool IsNonContiguous = false;
+ MapDimArrayTy Dims;
+ MapNonContiguousArrayTy Offsets;
+ MapNonContiguousArrayTy Counts;
+ MapNonContiguousArrayTy Strides;
+ };
+
private:
/// Kind that defines how a device pointer has to be returned.
struct MapInfo {
@@ -7222,9 +7236,11 @@
/// a flag marking the map as a pointer if requested. Add a flag marking the
/// map as the first one of a series of maps that relate to the same map
/// expression.
- OpenMPOffloadMappingFlags getMapTypeBits(
- OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
- bool IsImplicit, bool AddPtrFlag, bool AddIsTargetParamFlag) const {
+ OpenMPOffloadMappingFlags
+ getMapTypeBits(OpenMPMapClauseKind MapType,
+ ArrayRef<OpenMPMapModifierKind> MapModifiers, bool IsImplicit,
+ bool AddPtrFlag, bool AddIsTargetParamFlag,
+ bool IsNonContiguous) const {
OpenMPOffloadMappingFlags Bits =
IsImplicit ? OMP_MAP_IMPLICIT : OMP_MAP_NONE;
switch (MapType) {
@@ -7260,6 +7276,8 @@
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
!= MapModifiers.end())
Bits |= OMP_MAP_CLOSE;
+ if (IsNonContiguous)
+ Bits |= OMP_MAP_DESCRIPTOR;
return Bits;
}
@@ -7312,7 +7330,8 @@
MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
StructRangeInfoTy &PartialStruct, bool IsFirstComponentList,
- bool IsImplicit, bool ForDeviceAddr = false,
+ bool IsImplicit, StructNonContiguousInfo &NonContigInfo,
+ bool ForDeviceAddr = false,
ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
OverlappedElements = llvm::None) const {
// The following summarizes what has to be generated for each map and the
@@ -7561,6 +7580,12 @@
// whether we are dealing with a member of a declared struct.
const MemberExpr *EncounteredME = nullptr;
+ // Track for the total number of dimension. Start from one for the dummy
+ // dimension.
+ uint64_t DimSize = 1;
+
+ bool IsNonContiguous = NonContigInfo.IsNonContiguous;
+
for (; I != CE; ++I) {
// If the current component is member of a struct (parent struct) mark it.
if (!EncounteredME) {
@@ -7579,7 +7604,10 @@
// becomes the base address for the following components.
// A final array section, is one whose length can't be proved to be one.
+ // If the map item is non-contiguous then we don't treat any array section
+ // as final array section.
bool IsFinalArraySection =
+ !IsNonContiguous &&
isFinalArraySectionExpression(I->getAssociatedExpression());
// Get information on whether the element is a pointer. Have to do a
@@ -7597,7 +7625,10 @@
.getCanonicalType()
->isAnyPointerType()) ||
I->getAssociatedExpression()->getType()->isAnyPointerType();
- bool IsNonDerefPointer = IsPointer && !UO && !BO;
+ bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous;
+
+ if (OASE)
+ ++DimSize;
if (Next == CE || IsNonDerefPointer || IsFinalArraySection) {
// If this is not the last component, we expect the pointer to be
@@ -7653,7 +7684,7 @@
OMP_MAP_MEMBER_OF |
getMapTypeBits(MapType, MapModifiers, IsImplicit,
/*AddPtrFlag=*/false,
- /*AddIsTargetParamFlag=*/false);
+ /*AddIsTargetParamFlag=*/false, IsNonContiguous);
LB = BP;
llvm::Value *Size = nullptr;
// Do bitcopy of all non-overlapped structure elements.
@@ -7677,6 +7708,7 @@
Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty,
/*isSigned=*/true));
Types.push_back(Flags);
+ NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1);
LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
}
BasePointers.push_back(BP.getPointer());
@@ -7688,6 +7720,7 @@
Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
Types.push_back(Flags);
+ NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1);
break;
}
llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
@@ -7696,6 +7729,7 @@
Pointers.push_back(LB.getPointer());
Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
+ NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1);
// We need to add a pointer flag for each map that comes from the
// same expression except for the first one. We also need to signal
@@ -7704,7 +7738,7 @@
OpenMPOffloadMappingFlags Flags = getMapTypeBits(
MapType, MapModifiers, IsImplicit,
!IsExpressionFirstInfo || RequiresReference,
- IsCaptureFirstInfo && !RequiresReference);
+ IsCaptureFirstInfo && !RequiresReference, IsNonContiguous);
if (!IsExpressionFirstInfo) {
// If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
@@ -7764,6 +7798,180 @@
IsCaptureFirstInfo = false;
}
}
+
+ if (!IsNonContiguous)
+ return;
+
+ const ASTContext &Context = CGF.getContext();
+
+ // For supporting stride in array section, we need to initialize the first
+ // dimension size as 1, first offset as 0, and first count as 1
+ MapValuesArrayTy CurOffsets{llvm::ConstantInt::get(CGF.CGM.Int64Ty, 0)};
+ MapValuesArrayTy CurCounts{llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)};
+ MapValuesArrayTy CurStrides;
+ SmallVector<llvm::Value *, 4> DimSizes{
+ llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)};
+ uint64_t ElementTypeSize;
+
+ // Collect Size information for each dimension and get the element size as
+ // the first Stride. For example, for `int arr[10][10]`, the DimSizes
+ // should be [10, 10] and the first stride is 4 btyes.
+ for (const OMPClauseMappableExprCommon::MappableComponent &Component :
+ Components) {
+ const Expr *AssocExpr = Component.getAssociatedExpression();
+ const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+
+ if (!OASE)
+ continue;
+
+ QualType Ty = OMPArraySectionExpr::getBaseOriginalType(OASE->getBase());
+ auto *CAT = Context.getAsConstantArrayType(Ty);
+ auto *VAT = Context.getAsVariableArrayType(Ty);
+
+ // We need all the dimension size except for the last dimension.
+ assert((VAT || CAT || &Component == &*Components.begin()) &&
+ "Should be either ConstantArray or VariableArray if not the "
+ "first Component");
+
+ // Get element size if CurStrides is empty.
+ if (CurStrides.empty()) {
+ const Type *ElementType = nullptr;
+ if (CAT) {
+ 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.
+ if (ElementType) {
+ // For the case that having pointer as base, we need to remove one
+ // level of indirection.
+ if (&Component != &*Components.begin())
+ ElementType = ElementType->getPointeeOrArrayElementType();
+ ElementTypeSize =
+ Context.getTypeSizeInChars(ElementType).getQuantity();
+ CurStrides.push_back(
+ llvm::ConstantInt::get(CGF.Int64Ty, ElementTypeSize));
+ }
+ }
+ // Get dimension value except for the last dimension since we don't need
+ // it.
+ if (DimSizes.size() < Components.size() - 1) {
+ if (CAT) {
+ DimSizes.push_back(llvm::ConstantInt::get(
+ CGF.Int64Ty, CAT->getSize().getZExtValue()));
+ } else if (VAT) {
+ DimSizes.push_back(CGF.Builder.CreateIntCast(
+ CGF.EmitScalarExpr(VAT->getSizeExpr()), CGF.Int64Ty,
+ /*IsSigned=*/false));
+ }
+ }
+ }
+
+ // Skip the dummy dimension since we have already have its information.
+ auto DI = DimSizes.begin() + 1;
+ // Product of dimension.
+ llvm::Value *DimProd =
+ llvm::ConstantInt::get(CGF.CGM.Int64Ty, ElementTypeSize);
+
+ // Collect info for non-contiguous. Notice that offset, count, and stride
+ // are only meaningful for array-section, so we insert a null for anything
+ // other than array-section.
+ // Also, the size of offset, count, and stride are not the same as
+ // pointers, base_pointers, sizes, or dims. Instead, the size of offset,
+ // count, and stride are the same as the number of non-contiguous
+ // declaration in target update to/from clause.
+ for (const OMPClauseMappableExprCommon::MappableComponent &Component :
+ Components) {
+ const Expr *AssocExpr = Component.getAssociatedExpression();
+ AssocExpr->dump();
+ const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+ const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
+
+ if (AE) {
+ llvm::Value *Offset = CGF.Builder.CreateIntCast(
+ CGF.EmitScalarExpr(AE->getIdx()), CGF.Int64Ty,
+ /*isSigned=*/false);
+ CurOffsets.push_back(Offset);
+ CurCounts.push_back(llvm::ConstantInt::get(CGF.Int64Ty, /*V=*/1));
+ CurStrides.push_back(CurStrides.back());
+ continue;
+ }
+
+ if (!OASE)
+ continue;
+
+ // Offset
+ const Expr *OffsetExpr = OASE->getLowerBound();
+ llvm::Value *Offset = nullptr;
+ if (!OffsetExpr) {
+ // If offset is absent, then we just set it to zero.
+ Offset = llvm::ConstantInt::get(CGF.Int64Ty, 0);
+ } else {
+ Offset = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(OffsetExpr),
+ CGF.Int64Ty,
+ /*isSigned=*/false);
+ }
+ CurOffsets.push_back(Offset);
+
+ // Count
+ const Expr *CountExpr = OASE->getLength();
+ llvm::Value *Count = nullptr;
+ if (!CountExpr) {
+ // In Clang, once a high dimension is an array section, we construct all
+ // the lower dimension as array section, however, for case like
+ // arr[0:2][2], Clang construct the inner dimension as an array section
+ // but it actually is not in an array section form according to spec.
+ if (!OASE->getColonLocFirst().isValid() &&
+ !OASE->getColonLocSecond().isValid()) {
+ Count = llvm::ConstantInt::get(CGF.Int64Ty, 1);
+ } else {
+ // OpenMP 5.0, 2.1.5 Array Sections, Description.
+ // When the length is absent it defaults to ⌈(size −
+ // lower-bound)/stride⌉, where size is the size of the array
+ // dimension.
+ const Expr *StrideExpr = OASE->getStride();
+ llvm::Value *Stride =
+ StrideExpr
+ ? CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(StrideExpr),
+ CGF.Int64Ty, /*isSigned=*/false)
+ : llvm::ConstantInt::get(CGF.Int64Ty, /*V=*/1);
+ Count = CGF.Builder.CreateUDiv(CGF.Builder.CreateNUWSub(*DI, Offset),
+ Stride);
+ }
+ } else {
+ Count = CGF.EmitScalarExpr(CountExpr);
+ }
+ Count = CGF.Builder.CreateIntCast(Count, CGF.Int64Ty, /*isSigned=*/false);
+ CurCounts.push_back(Count);
+
+ // Stride_n' = Stride_n * (D_0 * D_1 ... * D_n-1) * Unit size
+ // Take `int arr[5][5][5]` and `arr[0:2:2][1:2:1][0:2:2]` as an example:
+ // Offset Count Stride
+ // D0 0 1 4 (int) <- dummy dimension
+ // D1 0 2 8 (2 * (1) * 4)
+ // D2 1 2 20 (1 * (1 * 5) * 4)
+ // D3 0 2 200 (2 * (1 * 5 * 4) * 4)
+ const Expr *StrideExpr = OASE->getStride();
+ llvm::Value *Stride =
+ StrideExpr
+ ? CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(StrideExpr),
+ CGF.Int64Ty, /*isSigned=*/false)
+ : llvm::ConstantInt::get(CGF.Int64Ty, 1);
+ DimProd = CGF.Builder.CreateNUWMul(DimProd, *(DI - 1));
+ CurStrides.push_back(CGF.Builder.CreateNUWMul(DimProd, Stride));
+ if (DI != DimSizes.end())
+ ++DI;
+ }
+
+ NonContigInfo.Offsets.push_back(CurOffsets);
+ NonContigInfo.Counts.push_back(CurCounts);
+ NonContigInfo.Strides.push_back(CurStrides);
}
/// Return the adjusted map modifiers if the declaration a capture refers to
@@ -7944,7 +8152,8 @@
/// index where it occurs is appended to the device pointers info array.
void generateAllInfo(MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
- MapFlagsArrayTy &Types) const {
+ MapFlagsArrayTy &Types,
+ CGOpenMPRuntime::TargetDataInfo &TargetDataInfo) const {
// We have to process the component lists that relate with the same
// declaration in a single chunk so that we can generate the map flags
// correctly. Therefore, we organize all lists in a map.
@@ -8120,6 +8329,7 @@
MapValuesArrayTy CurSizes;
MapFlagsArrayTy CurTypes;
StructRangeInfoTy PartialStruct;
+ StructNonContiguousInfo CurNonContigInfo;
for (const MapInfo &L : M.second) {
assert(!L.Components.empty() &&
@@ -8127,10 +8337,13 @@
// Remember the current base pointer index.
unsigned CurrentBasePointersIdx = CurBasePointers.size();
- generateInfoForComponentList(
- L.MapType, L.MapModifiers, L.Components, CurBasePointers,
- CurPointers, CurSizes, CurTypes, PartialStruct,
- IsFirstComponentList, L.IsImplicit, L.ForDeviceAddr);
+ CurNonContigInfo.IsNonContiguous =
+ L.Components.back().isNonContiguous();
+ generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
+ CurBasePointers, CurPointers, CurSizes,
+ CurTypes, PartialStruct,
+ IsFirstComponentList, L.IsImplicit,
+ CurNonContigInfo, L.ForDeviceAddr);
// If this entry relates with a device pointer, set the relevant
// declaration and add the 'return pointer' flag.
@@ -8193,6 +8406,14 @@
Pointers.append(CurPointers.begin(), CurPointers.end());
Sizes.append(CurSizes.begin(), CurSizes.end());
Types.append(CurTypes.begin(), CurTypes.end());
+ TargetDataInfo.Dims.append(CurNonContigInfo.Dims.begin(),
+ CurNonContigInfo.Dims.end());
+ TargetDataInfo.Offsets.append(CurNonContigInfo.Offsets.begin(),
+ CurNonContigInfo.Offsets.end());
+ TargetDataInfo.Counts.append(CurNonContigInfo.Counts.begin(),
+ CurNonContigInfo.Counts.end());
+ TargetDataInfo.Strides.append(CurNonContigInfo.Strides.begin(),
+ CurNonContigInfo.Strides.end());
}
}
@@ -8243,22 +8464,30 @@
MapValuesArrayTy CurSizes;
MapFlagsArrayTy CurTypes;
StructRangeInfoTy PartialStruct;
+ StructNonContiguousInfo CurNonContigInfo;
for (const MapInfo &L : M.second) {
assert(!L.Components.empty() &&
"Not expecting declaration with no component lists.");
- generateInfoForComponentList(
- L.MapType, L.MapModifiers, L.Components, CurBasePointers,
- CurPointers, CurSizes, CurTypes, PartialStruct,
- IsFirstComponentList, L.IsImplicit, L.ForDeviceAddr);
+ generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
+ CurBasePointers, CurPointers, CurSizes,
+ CurTypes, PartialStruct,
+ IsFirstComponentList, L.IsImplicit,
+ CurNonContigInfo, L.ForDeviceAddr);
IsFirstComponentList = false;
}
// 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())
+ if (PartialStruct.Base.isValid()) {
+ // Make sure Dims have the same size as BP, P, Sizes, and Types.
+ // Put 0 here to make sure that `emitOffloadingArrays` use it
+ // to skip processing this one. (OpenMP do not allow non-contigous for
+ // declare mapper)
+ CurNonContigInfo.Dims.push_back(0);
emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes,
PartialStruct);
+ }
// We need to append the results of this capture to what we already have.
BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
@@ -8503,6 +8732,7 @@
return *It == FD1;
});
}
+ StructNonContiguousInfo NonContigInfo;
// Associated with a capture, because the mapping flags depend on it.
// Go through all of the elements with the overlapped elements.
@@ -8518,7 +8748,7 @@
bool IsFirstComponentList = true;
generateInfoForComponentList(
MapType, MapModifiers, Components, BasePointers, Pointers, Sizes,
- Types, PartialStruct, IsFirstComponentList, IsImplicit,
+ Types, PartialStruct, IsFirstComponentList, IsImplicit, NonContigInfo,
/*ForDeviceAddr=*/false, OverlappedComponents);
}
// Go through other elements without overlapped elements.
@@ -8534,7 +8764,7 @@
generateInfoForComponentList(MapType, MapModifiers, Components,
BasePointers, Pointers, Sizes, Types,
PartialStruct, IsFirstComponentList,
- IsImplicit);
+ IsImplicit, NonContigInfo);
IsFirstComponentList = false;
}
}
@@ -8563,10 +8793,11 @@
!Res || *Res != OMPDeclareTargetDeclAttr::MT_Link)
continue;
StructRangeInfoTy PartialStruct;
+ StructNonContiguousInfo NonContigInfo;
generateInfoForComponentList(
C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers,
Pointers, Sizes, Types, PartialStruct,
- /*IsFirstComponentList=*/true, C->isImplicit());
+ /*IsFirstComponentList=*/true, C->isImplicit(), NonContigInfo);
assert(!PartialStruct.Base.isValid() &&
"No partial structs for declare target link expected.");
}
@@ -8659,6 +8890,78 @@
};
} // anonymous namespace
+static void emitNonContiguousDescriptor(
+ CodeGenFunction &CGF,
+ MappableExprsHandler::MapBaseValuesArrayTy &BasePointers,
+ MappableExprsHandler::MapValuesArrayTy &Pointers,
+ MappableExprsHandler::MapValuesArrayTy &Sizes,
+ MappableExprsHandler::MapFlagsArrayTy &MapTypes,
+ CGOpenMPRuntime::TargetDataInfo &Info) {
+ CodeGenModule &CGM = CGF.CGM;
+
+ // Build an array of struct descriptor_dim and then assign it to
+ // offload_args.
+ //
+ // struct descriptor_dim {
+ // uint64_t offset;
+ // uint64_t count;
+ // uint64_t stride
+ // };
+ ASTContext &C = CGF.getContext();
+ QualType Int64Ty = C.getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/true);
+ RecordDecl *RD;
+ RD = C.buildImplicitRecord("descriptor_dim");
+ RD->startDefinition();
+ addFieldToRecordDecl(C, RD, Int64Ty);
+ addFieldToRecordDecl(C, RD, Int64Ty);
+ addFieldToRecordDecl(C, RD, Int64Ty);
+ RD->completeDefinition();
+ QualType DimTy = C.getRecordType(RD);
+
+ enum { OffsetFD = 0, CountFD, StrideFD };
+ // The reason we need two index variable here is because the size of
+ // "Dims" is the same as the size of Components, however, the size of
+ // offset, count , and stride is equal to the size of base declaration
+ // that is non-contiguous.
+ for (unsigned I = 0, L = 0, E = Info.Dims.size(); I < E; ++I) {
+ // Skip emitting ir if dimension is 1.
+ if (Info.Dims[I] == 1)
+ continue;
+ // For supporting stride, each descriptor has a dummy dimension, hence,
+ // adding one to the original dimension size.
+ llvm::APInt Size(/*numBits=*/32, Info.Dims[I]);
+ QualType ArrayTy =
+ C.getConstantArrayType(DimTy, Size, nullptr, ArrayType::Normal, 0);
+ Address DimsAddr = CGF.CreateMemTemp(ArrayTy, "dims");
+ for (unsigned II = 0, EE = Info.Dims[I]; II < EE; ++II) {
+ unsigned RevIdx = EE - II - 1;
+ LValue DimsLVal = CGF.MakeAddrLValue(
+ CGF.Builder.CreateConstArrayGEP(DimsAddr, II), DimTy);
+ // Offset
+ LValue OffsetLVal = CGF.EmitLValueForField(
+ DimsLVal, *std::next(RD->field_begin(), OffsetFD));
+ CGF.EmitStoreOfScalar(Info.Offsets[L][RevIdx], OffsetLVal);
+ // Count
+ LValue CountLVal = CGF.EmitLValueForField(
+ DimsLVal, *std::next(RD->field_begin(), CountFD));
+ CGF.EmitStoreOfScalar(Info.Counts[L][RevIdx], CountLVal);
+ // Stride
+ LValue StrideLVal = CGF.EmitLValueForField(
+ DimsLVal, *std::next(RD->field_begin(), StrideFD));
+ CGF.EmitStoreOfScalar(Info.Strides[L][RevIdx], StrideLVal);
+ }
+ // args[I] = &dims
+ Address DAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+ DimsAddr, CGM.Int8PtrTy);
+ llvm::Value *P = CGF.Builder.CreateConstInBoundsGEP2_32(
+ llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
+ Info.PointersArray, 0, I);
+ Address PAddr(P, CGF.getPointerAlign());
+ CGF.Builder.CreateStore(DAddr.getPointer(), PAddr);
+ ++L;
+ }
+}
+
/// Emit the arrays used to pass the captures and map information to the
/// offloading runtime library. If there is no map or capture information,
/// return nullptr by reference.
@@ -8668,7 +8971,8 @@
MappableExprsHandler::MapValuesArrayTy &Pointers,
MappableExprsHandler::MapValuesArrayTy &Sizes,
MappableExprsHandler::MapFlagsArrayTy &MapTypes,
- CGOpenMPRuntime::TargetDataInfo &Info) {
+ CGOpenMPRuntime::TargetDataInfo &Info,
+ bool IsNonContiguous = false) {
CodeGenModule &CGM = CGF.CGM;
ASTContext &Ctx = CGF.getContext();
@@ -8711,8 +9015,15 @@
// We expect all the sizes to be constant, so we collect them to create
// a constant array.
SmallVector<llvm::Constant *, 16> ConstSizes;
- for (llvm::Value *S : Sizes)
- ConstSizes.push_back(cast<llvm::Constant>(S));
+ for (unsigned I = 0, E = Sizes.size(); I < E; ++I) {
+ if (IsNonContiguous &&
+ (MapTypes[I] & MappableExprsHandler::OMP_MAP_DESCRIPTOR)) {
+ ConstSizes.push_back(
+ llvm::ConstantInt::get(CGF.Int64Ty, Info.Dims[I]));
+ } else {
+ ConstSizes.push_back(cast<llvm::Constant>(Sizes[I]));
+ }
+ }
auto *SizesArrayInit = llvm::ConstantArray::get(
llvm::ArrayType::get(CGM.Int64Ty, ConstSizes.size()), ConstSizes);
@@ -8776,6 +9087,12 @@
}
}
}
+
+ if (!IsNonContiguous || Info.Offsets.empty() || Info.NumberOfPtrs == 0)
+ return;
+
+ emitNonContiguousDescriptor(CGF, BasePointers, Pointers, Sizes, MapTypes,
+ Info);
}
/// Emit the arguments to be passed to the runtime library based on the
@@ -10138,10 +10455,11 @@
// Get map clause information.
MappableExprsHandler MCHandler(D, CGF);
- MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+ MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Info);
// Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info,
+ /*IsNonContiguous=*/true);
llvm::Value *BasePointersArrayArg = nullptr;
llvm::Value *PointersArrayArg = nullptr;
@@ -10379,11 +10697,12 @@
// Get map clause information.
MappableExprsHandler MEHandler(D, CGF);
- MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
-
TargetDataInfo Info;
+ MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Info);
+
// Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info,
+ /*IsNonContiguous=*/true);
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
Info.PointersArray, Info.SizesArray,
Info.MapTypesArray, Info);
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9778,6 +9778,8 @@
"bit fields cannot be used to specify storage in a '%0' clause">;
def err_array_section_does_not_specify_contiguous_storage : Error<
"array section does not specify contiguous storage">;
+def err_array_section_does_not_specify_length : Error<
+ "array section does not specify length for outermost dimension">;
def err_omp_union_type_not_allowed : Error<
"mapping of union members is not allowed">;
def err_omp_expected_access_to_data_field : Error<
Index: clang/include/clang/AST/OpenMPClause.h
===================================================================
--- clang/include/clang/AST/OpenMPClause.h
+++ clang/include/clang/AST/OpenMPClause.h
@@ -27,6 +27,7 @@
#include "clang/Basic/SourceLocation.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/MapVector.h"
+#include "llvm/ADT/PointerIntPair.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/iterator.h"
#include "llvm/ADT/iterator_range.h"
@@ -4737,8 +4738,9 @@
/// subscript it may not have any associated declaration. In that case the
/// associated declaration is set to nullptr.
class MappableComponent {
- /// Expression associated with the component.
- Expr *AssociatedExpression = nullptr;
+ /// Pair of Expression and Non-contiguous pair associated with the
+ /// component.
+ llvm::PointerIntPair<Expr *, 1, bool> AssociatedExpressionNonContiguousPr;
/// Declaration associated with the declaration. If the component does
/// not have a declaration (e.g. array subscripts or section), this is set
@@ -4748,14 +4750,22 @@
public:
explicit MappableComponent() = default;
explicit MappableComponent(Expr *AssociatedExpression,
- ValueDecl *AssociatedDeclaration)
- : AssociatedExpression(AssociatedExpression),
+ ValueDecl *AssociatedDeclaration,
+ bool IsNonContiguous)
+ : AssociatedExpressionNonContiguousPr(AssociatedExpression,
+ IsNonContiguous),
AssociatedDeclaration(
AssociatedDeclaration
? cast<ValueDecl>(AssociatedDeclaration->getCanonicalDecl())
: nullptr) {}
- Expr *getAssociatedExpression() const { return AssociatedExpression; }
+ Expr *getAssociatedExpression() const {
+ return AssociatedExpressionNonContiguousPr.getPointer();
+ }
+
+ bool isNonContiguous() const {
+ return AssociatedExpressionNonContiguousPr.getInt();
+ }
ValueDecl *getAssociatedDeclaration() const {
return AssociatedDeclaration;
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
