cchen updated this revision to Diff 273517.
cchen added a comment.
Fix based on feedback
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D79972/new/
https://reviews.llvm.org/D79972
Files:
clang/include/clang/AST/OpenMPClause.h
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 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} expected-error {{invalid operands to binary expression ('S8 *' and 'S8 *')}}
{}
+
+ double marr[10][5][10];
+#pragma omp target update to(marr [0:] [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,6 +1,8 @@
-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -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}}
@@ -36,5 +38,21 @@
{
foo();
}
+
+ double marr[10][5][10];
+#pragma omp target update to(marr [0:] [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: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'}}
+
+ 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'}}
+
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,304 @@
#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 3, i64 4, i64 2]
+// 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 [3 x [[STRUCT_DESCRIPTOR]]],
+ // CK19: [[DIMS_2:%.+]] = alloca [2 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: [[LEN:%.+]] = sub nuw i64 4, [[ARG_ADDR:%.+]]
+ // 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 [3 x [[STRUCT_DESCRIPTOR]]], [3 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 [3 x [[STRUCT_DESCRIPTOR]]], [3 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 [3 x [[STRUCT_DESCRIPTOR]]], [3 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: [[PC0:%.+]] = bitcast [3 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_4:%.+]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR]]], [2 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 0
+ // CK19: [[OFFSET_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 0, i64* [[OFFSET_2_1]],
+ // CK19: [[COUNT_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 2, i64* [[COUNT_2_1]],
+ // CK19: [[STRIDE_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 12, i64* [[STRIDE_2_1]],
+ // CK19: [[DIM_5:%.+]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR]]], [2 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 1
+ // CK19: [[OFFSET_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 0
+ // CK19: store i64 1, i64* [[OFFSET_2_2]],
+ // CK19: [[COUNT_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 1
+ // CK19: store i64 2, i64* [[COUNT_2_2]],
+ // CK19: [[STRIDE_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 2
+ // CK19: store i64 4, i64* [[STRIDE_2_2]],
+ // CK19: [[PC1:%.+]] = bitcast [2 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 2]
+// CK20: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK20-LABEL: _Z3foo
+void foo(int arg) {
+ ST arr[3][4];
+ // CK20: [[DIMS:%.+]] = alloca [2 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 [2 x [[STRUCT_DESCRIPTOR]]], [2 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 [2 x [[STRUCT_DESCRIPTOR]]], [2 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-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 [2 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 [3 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 [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+ // CK21: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+ // CK21: store i64 0, i64* [[OFFSET]],
+ // CK21: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+ // CK21: store i64 2, i64* [[COUNT]],
+ // CK21: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+ // CK21: store i64 {{400|800}}, i64* [[STRIDE]],
+ // CK21: [[DIM_2:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+ // CK21: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+ // CK21: store i64 1, i64* [[OFFSET]],
+ // CK21: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+ // CK21: store i64 3, i64* [[COUNT]],
+ // CK21: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+ // CK21: store i64 {{40|80}}, i64* [[STRIDE]],
+ // CK21: [[DIM_3:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+ // CK21: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+ // CK21: store i64 0, i64* [[OFFSET]],
+ // CK21: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+ // CK21: store i64 4, i64* [[COUNT]],
+ // CK21: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+ // CK21: store i64 {{4|8}}, i64* [[STRIDE]],
+ // 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 [3 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 3]
+// CK22: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+struct ST {
+ // CK22: _ZN2ST3fooEPA10_Pi
+ void foo(int *arr[5][10]) {
+ // CK22: [[DIMS:%.+]] = alloca [3 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 [3 x [[STRUCT_DESCRIPTOR]]], [3 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 [3 x [[STRUCT_DESCRIPTOR]]], [3 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 [3 x [[STRUCT_DESCRIPTOR]]], [3 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-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 [3 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
#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
@@ -20,23 +28,64 @@
#pragma omp target update to(([a][targ])p, a) if(l>5) device(l) nowait depend(inout:l)
#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 to(marr[:] [0:2] [0:2])
+
+#pragma omp target update from(marr[:] [0:2] [0:2])
+
+#pragma omp target update to(marr[:][:l] [l:])
+
+#pragma omp target update from(marr[:][: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])
+
+// 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 to(marr[:][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:][:l][l:])
+// OMP5-NEXT: #pragma omp target update from(marr[:][: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])
+#endif
+
return a + targ + (T)b;
}
// CHECK: static T a, *p;
// CHECK-NEXT: U b;
-// 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: static int a, *p;
-// CHECK-NEXT: float b;
-// 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: static char a, *p;
-// CHECK-NEXT: float b;
-// 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)
int main(int argc, char **argv) {
static int a;
@@ -50,6 +99,40 @@
// 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[:] [0:2] [0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:][0:2][0:2])
+#pragma omp target update from(marr[:] [0:2] [0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:][0:2][0:2])
+#pragma omp target update to(marr[:][:n] [n:])
+// OMP5: #pragma omp target update to(marr[:][: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])
+#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
@@ -6581,6 +6581,7 @@
Record.push_back(N);
for (auto &M : C->all_components()) {
Record.AddStmt(M.getAssociatedExpression());
+ Record.writeBool(M.isNonContiguous());
Record.AddDeclRef(M.getAssociatedDeclaration());
}
}
@@ -6605,6 +6606,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
@@ -12510,10 +12510,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);
}
@@ -12627,10 +12627,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);
}
@@ -12677,10 +12677,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);
}
@@ -12727,10 +12727,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);
}
@@ -12771,8 +12771,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);
}
@@ -12814,8 +12814,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.
@@ -3519,6 +3519,7 @@
if (isOpenMPTargetExecutionDirective(DKind)) {
OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map,
+ Stack->getCurrentDirective(),
/*NoDiagnose=*/true))
return;
const auto *VD = cast<ValueDecl>(
@@ -16436,11 +16437,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;
@@ -16465,7 +16469,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;
}
@@ -16537,7 +16541,7 @@
AllowWholeSizeArraySection = false;
// Record the component.
- Components.emplace_back(ME, FD);
+ Components.emplace_back(ME, FD, IsNonContiguous);
return RelevantExpr || Visit(E);
}
@@ -16575,7 +16579,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);
}
@@ -16614,6 +16618,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.
@@ -16623,6 +16634,9 @@
return false;
}
+ if (IsPointer)
+ AllowAnotherPtr = false;
+
if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
Expr::EvalResult ResultR;
Expr::EvalResult ResultL;
@@ -16648,14 +16662,14 @@
}
// Record the component - we don't have any declaration associated.
- Components.emplace_back(OASE, nullptr);
+ Components.emplace_back(OASE, nullptr, 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());
}
@@ -16668,7 +16682,7 @@
}
if (!RelevantExpr) {
// Record the component if haven't found base decl.
- Components.emplace_back(UO, nullptr);
+ Components.emplace_back(UO, nullptr, false);
}
return RelevantExpr || Visit(UO->getSubExpr()->IgnoreParenImpCasts());
}
@@ -16684,7 +16698,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");
@@ -16695,7 +16709,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 *) {
@@ -16706,10 +16720,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
@@ -16721,10 +16735,10 @@
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()))
return Checker.getFoundBase();
@@ -17204,7 +17218,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;
@@ -18452,8 +18467,10 @@
// 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(
+ OMPClauseMappableExprCommon::MappableComponent(
+ SimpleRefExpr, D,
+ /*IsNonContiguous=*/false));
}
if (MVLI.ProcessedVarList.empty())
@@ -18500,8 +18517,8 @@
// only need a component.
MVLI.VarBaseDeclarations.push_back(D);
MVLI.VarComponents.emplace_back();
- MVLI.VarComponents.back().push_back(
- OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D));
+ MVLI.VarComponents.back().emplace_back(SimpleRefExpr, D,
+ /*IsNonContiguous=*/false);
}
if (MVLI.ProcessedVarList.empty())
@@ -18567,7 +18584,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
@@ -1619,6 +1619,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
@@ -7043,6 +7043,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,
@@ -7078,6 +7082,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.
@@ -7091,6 +7097,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 {
@@ -7221,9 +7235,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) {
@@ -7259,6 +7275,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,7 @@
MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
StructRangeInfoTy &PartialStruct, bool IsFirstComponentList,
- bool IsImplicit,
+ bool IsImplicit, StructNonContiguousInfo &NonContigInfo,
ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
OverlappedElements = llvm::None) const {
// The following summarizes what has to be generated for each map and the
@@ -7561,6 +7579,11 @@
// whether we are dealing with a member of a declared struct.
const MemberExpr *EncounteredME = nullptr;
+ // Track for the total number of dimension.
+ uint64_t DimSize = 0;
+
+ 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 +7602,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 +7623,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 +7682,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 +7706,7 @@
Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty,
/*isSigned=*/true));
Types.push_back(Flags);
+ NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 0);
LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
}
BasePointers.push_back(BP.getPointer());
@@ -7688,6 +7718,7 @@
Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
Types.push_back(Flags);
+ NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 0);
break;
}
llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
@@ -7696,6 +7727,7 @@
Pointers.push_back(LB.getPointer());
Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
+ NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 0);
// 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 +7736,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 +7796,141 @@
IsCaptureFirstInfo = false;
}
}
+
+ if (IsNonContiguous) {
+ const ASTContext &Context = CGF.getContext();
+
+ MapValuesArrayTy CurOffsets;
+ MapValuesArrayTy CurCounts;
+ MapValuesArrayTy CurStrides;
+ llvm::Value *CurStride = nullptr;
+ SmallVector<llvm::Value *, 4> DimSizes;
+
+ // 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;
+ 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;
+ uint64_t ElementTypeSize;
+ 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));
+ }
+ }
+ }
+
+ // We need dimension size to compute stride
+ auto DI = DimSizes.begin();
+
+ // 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();
+ const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+
+ if (!OASE)
+ continue;
+
+ // Offset
+ const Expr *OffsetExpr = nullptr;
+ 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) {
+ // If length is absent then we calculate it as (Total length -
+ // lower_bound)
+ Count = CGF.Builder.CreateNUWSub(*DI, Offset);
+ } else {
+ Count = CGF.EmitScalarExpr(CountExpr);
+ }
+ Count =
+ CGF.Builder.CreateIntCast(Count, CGF.Int64Ty, /*isSigned=*/false);
+ CurCounts.push_back(Count);
+
+ // Stride = previous stride * previous dimension size
+ // Take `int arr[5][10]` and `arr[0:2][0:2]` as an example:
+ // Dimension 1 Dimension 0
+ // Offset 0 0
+ // Count 2 2
+ // Stride 40 bytes (4x10) 4 bytes (int)
+ if (DI != DimSizes.end()) {
+ CurStride = CGF.Builder.CreateNUWMul(CurStrides.back(), *DI);
+ CurStrides.push_back(CurStride);
+ ++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 +8111,10 @@
/// 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, MapDimArrayTy &Dims,
+ MapNonContiguousArrayTy &Offsets,
+ MapNonContiguousArrayTy &Counts,
+ MapNonContiguousArrayTy &Strides) 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.
@@ -8055,6 +8225,7 @@
MapValuesArrayTy CurSizes;
MapFlagsArrayTy CurTypes;
StructRangeInfoTy PartialStruct;
+ StructNonContiguousInfo CurNonContigInfo;
for (const MapInfo &L : M.second) {
assert(!L.Components.empty() &&
@@ -8062,10 +8233,12 @@
// 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);
+ CurNonContigInfo.IsNonContiguous =
+ L.Components.back().isNonContiguous();
+ generateInfoForComponentList(
+ L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+ CurPointers, CurSizes, CurTypes, PartialStruct,
+ IsFirstComponentList, L.IsImplicit, CurNonContigInfo);
// If this entry relates with a device pointer, set the relevant
// declaration and add the 'return pointer' flag.
@@ -8114,6 +8287,13 @@
Pointers.append(CurPointers.begin(), CurPointers.end());
Sizes.append(CurSizes.begin(), CurSizes.end());
Types.append(CurTypes.begin(), CurTypes.end());
+ Dims.append(CurNonContigInfo.Dims.begin(), CurNonContigInfo.Dims.end());
+ Offsets.append(CurNonContigInfo.Offsets.begin(),
+ CurNonContigInfo.Offsets.end());
+ Counts.append(CurNonContigInfo.Counts.begin(),
+ CurNonContigInfo.Counts.end());
+ Strides.append(CurNonContigInfo.Strides.begin(),
+ CurNonContigInfo.Strides.end());
}
}
@@ -8164,22 +8344,29 @@
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);
+ generateInfoForComponentList(
+ L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+ CurPointers, CurSizes, CurTypes, PartialStruct,
+ IsFirstComponentList, L.IsImplicit, CurNonContigInfo);
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());
@@ -8424,6 +8611,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.
@@ -8437,10 +8625,10 @@
ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
OverlappedComponents = Pair.getSecond();
bool IsFirstComponentList = true;
- generateInfoForComponentList(MapType, MapModifiers, Components,
- BasePointers, Pointers, Sizes, Types,
- PartialStruct, IsFirstComponentList,
- IsImplicit, OverlappedComponents);
+ generateInfoForComponentList(
+ MapType, MapModifiers, Components, BasePointers, Pointers, Sizes,
+ Types, PartialStruct, IsFirstComponentList, IsImplicit, NonContigInfo,
+ OverlappedComponents);
}
// Go through other elements without overlapped elements.
bool IsFirstComponentList = OverlappedData.empty();
@@ -8455,7 +8643,7 @@
generateInfoForComponentList(MapType, MapModifiers, Components,
BasePointers, Pointers, Sizes, Types,
PartialStruct, IsFirstComponentList,
- IsImplicit);
+ IsImplicit, NonContigInfo);
IsFirstComponentList = false;
}
}
@@ -8484,10 +8672,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.");
}
@@ -8580,6 +8769,75 @@
};
} // 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) {
+ if (Info.Dims[I] == 0)
+ continue;
+ 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.
@@ -8589,7 +8847,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();
@@ -8632,8 +8891,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);
@@ -8697,6 +8963,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
@@ -10058,10 +10330,13 @@
// Get map clause information.
MappableExprsHandler MCHandler(D, CGF);
- MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+ MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes,
+ Info.Dims, Info.Offsets, Info.Counts,
+ Info.Strides);
// 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;
@@ -10298,11 +10573,14 @@
// Get map clause information.
MappableExprsHandler MEHandler(D, CGF);
- MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
-
TargetDataInfo Info;
+ MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes,
+ Info.Dims, Info.Offsets, Info.Counts,
+ Info.Strides);
+
// 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/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
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
