cchen created this revision.
cchen added a reviewer: ABataev.
Herald added subscribers: cfe-commits, 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.
More details can be found here:
https://github.com/chichunchen/openmp-50-design/blob/master/target_update_noncontiguous.pptx
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D79972
Files:
clang/include/clang/AST/OpenMPClause.h
clang/lib/AST/OpenMPClause.cpp
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
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,142 @@
#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 [1 x i64] [i64 3]
+// CK19: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 2081]
+
+// CK19-LABEL: _Z3foo
+void foo(int arg) {
+ int arr[3][4][5];
+
+ // CK19: [[DIMS:%.+]] = 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: [[LEN:%.+]] = sub nuw i64 4, [[ARG_ADDR:%.+]]
+ // CK19: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], i{{.+}} 0, i{{.+}} 0
+ // CK19: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 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-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]]{{.+}})
+ // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK19-DAG: [[PC0:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+ // CK19-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+ // CK19-DAG: store i8* [[PC0]], i8** [[PTRS]],
+
+#pragma omp target update to(arr[0:2][arg:][1:4])
+ {++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 2081]
+
+// 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
#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
@@ -6511,6 +6511,8 @@
Record.AddStmt(M.getAssociatedExpression());
Record.AddDeclRef(M.getAssociatedDeclaration());
}
+ for (auto NC : C->non_contiguous_lists())
+ Record.push_back(NC);
}
void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
@@ -6535,6 +6537,8 @@
Record.AddStmt(M.getAssociatedExpression());
Record.AddDeclRef(M.getAssociatedDeclaration());
}
+ for (auto NC : C->non_contiguous_lists())
+ Record.push_back(NC);
}
void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
Index: clang/lib/Serialization/ASTReader.cpp
===================================================================
--- clang/lib/Serialization/ASTReader.cpp
+++ clang/lib/Serialization/ASTReader.cpp
@@ -12513,6 +12513,13 @@
AssociatedExpr, AssociatedDecl));
}
C->setComponents(Components, ListSizes);
+
+ SmallVector<bool, 16> ListNonContiguous;
+ ListNonContiguous.reserve(TotalLists);
+ for (unsigned i = 0; i < TotalLists; ++i) {
+ ListNonContiguous.push_back(Record.readBool());
+ }
+ C->setNonContiguousLists(ListNonContiguous);
}
void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
@@ -12563,6 +12570,13 @@
AssociatedExpr, AssociatedDecl));
}
C->setComponents(Components, ListSizes);
+
+ SmallVector<bool, 16> ListNonContiguous;
+ ListNonContiguous.reserve(TotalLists);
+ for (unsigned i = 0; i < TotalLists; ++i) {
+ ListNonContiguous.push_back(Record.readBool());
+ }
+ C->setNonContiguousLists(ListNonContiguous);
}
void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -47,7 +47,8 @@
static const Expr *checkMapClauseExpressionBase(
Sema &SemaRef, Expr *E,
OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
- OpenMPClauseKind CKind, bool NoDiagnose);
+ bool &IsNonContiguous, OpenMPClauseKind CKind, OpenMPDirectiveKind DKind,
+ bool NoDiagnose);
namespace {
/// Default data sharing attributes, which can be applied to directive.
@@ -3395,7 +3396,10 @@
}
if (isOpenMPTargetExecutionDirective(DKind)) {
OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
- if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map,
+ bool IsNonContiguous = false;
+ if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents,
+ IsNonContiguous, OMPC_map,
+ Stack->getCurrentDirective(),
/*NoDiagnose=*/true))
return;
const auto *VD = cast<ValueDecl>(
@@ -16142,7 +16146,9 @@
class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
Sema &SemaRef;
OpenMPClauseKind CKind = OMPC_unknown;
+ OpenMPDirectiveKind DKind = OMPD_unknown;
OMPClauseMappableExprCommon::MappableExprComponentList &Components;
+ bool &IsNonContiguousRef;
bool NoDiagnose = false;
const Expr *RelevantExpr = nullptr;
bool AllowUnitySizeArraySection = true;
@@ -16320,6 +16326,9 @@
// pointer. Otherwise, only unitary sections are accepted.
if (NotWhole || IsPointer)
AllowWholeSizeArraySection = false;
+ } else if (DKind == OMPD_target_update &&
+ SemaRef.getLangOpts().OpenMP >= 50) {
+ IsNonContiguousRef = 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.
@@ -16412,11 +16421,13 @@
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),
- NoDiagnose(NoDiagnose), ELoc(ELoc), ERange(ERange) {}
+ bool &IsNonContiguousTargetUpdate, bool NoDiagnose, SourceLocation &ELoc,
+ SourceRange &ERange)
+ : SemaRef(SemaRef), CKind(CKind), DKind(DKind), Components(Components),
+ IsNonContiguousRef(IsNonContiguousTargetUpdate), NoDiagnose(NoDiagnose),
+ ELoc(ELoc), ERange(ERange) {}
};
} // namespace
@@ -16427,11 +16438,12 @@
static const Expr *checkMapClauseExpressionBase(
Sema &SemaRef, Expr *E,
OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
- OpenMPClauseKind CKind, bool NoDiagnose) {
+ bool &IsNonContiguousTargetUpdate, OpenMPClauseKind CKind,
+ OpenMPDirectiveKind DKind, bool NoDiagnose) {
SourceLocation ELoc = E->getExprLoc();
SourceRange ERange = E->getSourceRange();
- MapBaseChecker Checker(SemaRef, CKind, CurComponents, NoDiagnose, ELoc,
- ERange);
+ MapBaseChecker Checker(SemaRef, CKind, DKind, CurComponents,
+ IsNonContiguousTargetUpdate, NoDiagnose, ELoc, ERange);
if (Checker.Visit(E->IgnoreParens()))
return Checker.getFoundBase();
return nullptr;
@@ -16809,6 +16821,8 @@
SmallVector<ValueDecl *, 16> VarBaseDeclarations;
// The reference to the user-defined mapper associated with every expression.
SmallVector<Expr *, 16> UDMapperList;
+ // The list of whether the expression is non-contiguous or not
+ SmallVector<bool, 16> IsNonContiguousList;
MappableVarListInfo(ArrayRef<Expr *> VarList) : VarList(VarList) {
// We have a list of components and base declarations for each entry in the
@@ -16905,12 +16919,15 @@
}
OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
+ bool IsNonContiguousTargetUpdate = false;
ValueDecl *CurDeclaration = nullptr;
// 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, IsNonContiguousTargetUpdate, CKind,
+ DSAS->getCurrentDirective(),
+ /*NoDiagnose=*/false);
if (!BE)
continue;
@@ -16933,6 +16950,7 @@
MVLI.VarComponents.back().append(CurComponents.begin(),
CurComponents.end());
MVLI.VarBaseDeclarations.push_back(nullptr);
+ MVLI.IsNonContiguousList.push_back(IsNonContiguousTargetUpdate);
continue;
}
@@ -17110,6 +17128,7 @@
CurComponents.end());
MVLI.VarBaseDeclarations.push_back(isa<MemberExpr>(BE) ? nullptr
: CurDeclaration);
+ MVLI.IsNonContiguousList.push_back(IsNonContiguousTargetUpdate);
}
}
@@ -17147,11 +17166,11 @@
// We need to produce a map clause even if we don't have variables so that
// other diagnostics related with non-existing map clauses are accurate.
- return OMPMapClause::Create(Context, Locs, MVLI.ProcessedVarList,
- MVLI.VarBaseDeclarations, MVLI.VarComponents,
- MVLI.UDMapperList, Modifiers, ModifiersLoc,
- MapperIdScopeSpec.getWithLocInContext(Context),
- MapperId, MapType, IsMapTypeImplicit, MapLoc);
+ return OMPMapClause::Create(
+ Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
+ MVLI.VarComponents, MVLI.IsNonContiguousList, MVLI.UDMapperList,
+ Modifiers, ModifiersLoc, MapperIdScopeSpec.getWithLocInContext(Context),
+ MapperId, MapType, IsMapTypeImplicit, MapLoc);
}
QualType Sema::ActOnOpenMPDeclareReductionType(SourceLocation TyLoc,
@@ -18063,7 +18082,7 @@
return OMPToClause::Create(
Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
- MVLI.VarComponents, MVLI.UDMapperList,
+ MVLI.VarComponents, MVLI.IsNonContiguousList, MVLI.UDMapperList,
MapperIdScopeSpec.getWithLocInContext(Context), MapperId);
}
@@ -18080,7 +18099,7 @@
return OMPFromClause::Create(
Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
- MVLI.VarComponents, MVLI.UDMapperList,
+ MVLI.VarComponents, MVLI.IsNonContiguousList, MVLI.UDMapperList,
MapperIdScopeSpec.getWithLocInContext(Context), MapperId);
}
Index: clang/lib/CodeGen/CGOpenMPRuntime.h
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.h
+++ clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -1577,6 +1577,10 @@
llvm::Value *SizesArray = nullptr;
/// The array of map types passed to the runtime library.
llvm::Value *MapTypesArray = nullptr;
+ /// The array of array of dims passed to the runtime library.
+ llvm::Value *DimsArray = nullptr;
+ /// The array of array of descriptor passed to the runtime library.
+ llvm::Value *DescriptorsArray = nullptr;
/// The total number of pointers passed to the runtime library.
unsigned NumberOfPtrs = 0u;
/// Map between the a declaration of a capture and the corresponding base
@@ -1592,6 +1596,8 @@
PointersArray = nullptr;
SizesArray = nullptr;
MapTypesArray = nullptr;
+ DimsArray = nullptr;
+ DescriptorsArray = nullptr;
NumberOfPtrs = 0u;
}
/// Return true if the current target data information has valid arrays.
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7643,6 +7643,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 = 0x800,
/// The 16 MSBs of the flags indicate whether the entry is member of some
/// struct/class.
OMP_MAP_MEMBER_OF = 0xffff000000000000,
@@ -7678,6 +7682,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.
@@ -7699,15 +7705,17 @@
ArrayRef<OpenMPMapModifierKind> MapModifiers;
bool ReturnDevicePointer = false;
bool IsImplicit = false;
+ bool IsNonContiguous = false;
MapInfo() = default;
MapInfo(
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
OpenMPMapClauseKind MapType,
- ArrayRef<OpenMPMapModifierKind> MapModifiers,
- bool ReturnDevicePointer, bool IsImplicit)
+ ArrayRef<OpenMPMapModifierKind> MapModifiers, bool ReturnDevicePointer,
+ bool IsImplicit, bool IsNonContiguous)
: Components(Components), MapType(MapType), MapModifiers(MapModifiers),
- ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {}
+ ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit),
+ IsNonContiguous(IsNonContiguous) {}
};
/// If use_device_ptr is used on a pointer which is a struct member and there
@@ -7821,9 +7829,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) {
@@ -7859,6 +7869,8 @@
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
!= MapModifiers.end())
Bits |= OMP_MAP_CLOSE;
+ if (IsNonContiguous)
+ Bits |= OMP_MAP_DESCRIPTOR;
return Bits;
}
@@ -7910,11 +7922,11 @@
ArrayRef<OpenMPMapModifierKind> MapModifiers,
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
- MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
+ MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, MapDimArrayTy &Dims,
StructRangeInfoTy &PartialStruct, bool IsFirstComponentList,
bool IsImplicit,
ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
- OverlappedElements = llvm::None) const {
+ OverlappedElements = llvm::None, bool IsNonContiguous = false) const {
// The following summarizes what has to be generated for each map and the
// types below. The generated information is expressed in this order:
// base pointer, section pointer, size, flags
@@ -8161,6 +8173,9 @@
// 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;
+
for (; I != CE; ++I) {
// If the current component is member of a struct (parent struct) mark it.
if (!EncounteredME) {
@@ -8179,8 +8194,11 @@
// 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 =
- isFinalArraySectionExpression(I->getAssociatedExpression());
+ isFinalArraySectionExpression(I->getAssociatedExpression()) &&
+ (!IsNonContiguous);
// Get information on whether the element is a pointer. Have to do a
// special treatment for array sections given that they are built-in
@@ -8199,6 +8217,11 @@
I->getAssociatedExpression()->getType()->isAnyPointerType();
bool IsNonDerefPointer = IsPointer && !UO && !BO;
+ if (OASE || OAShE ||
+ dyn_cast<ArraySubscriptExpr>(I->getAssociatedExpression())) {
+ DimSize++;
+ }
+
if (Next == CE || IsNonDerefPointer || IsFinalArraySection) {
// If this is not the last component, we expect the pointer to be
// associated with an array expression or member expression.
@@ -8253,7 +8276,8 @@
OMP_MAP_MEMBER_OF |
getMapTypeBits(MapType, MapModifiers, IsImplicit,
/*AddPtrFlag=*/false,
- /*AddIsTargetParamFlag=*/false);
+ /*AddIsTargetParamFlag=*/false,
+ /*IsNonContiguous=*/IsNonContiguous);
LB = BP;
llvm::Value *Size = nullptr;
// Do bitcopy of all non-overlapped structure elements.
@@ -8277,6 +8301,7 @@
Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty,
/*isSigned=*/true));
Types.push_back(Flags);
+ Dims.push_back(IsNonContiguous ? DimSize : 0);
LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
}
BasePointers.push_back(BP.getPointer());
@@ -8288,6 +8313,7 @@
Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
Types.push_back(Flags);
+ Dims.push_back(IsNonContiguous ? DimSize : 0);
break;
}
llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
@@ -8296,15 +8322,17 @@
Pointers.push_back(LB.getPointer());
Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
+ 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
// this map is the first one that relates with the current capture
// (there is a set of entries for each capture).
- OpenMPOffloadMappingFlags Flags = getMapTypeBits(
- MapType, MapModifiers, IsImplicit,
- !IsExpressionFirstInfo || RequiresReference,
- IsCaptureFirstInfo && !RequiresReference);
+ OpenMPOffloadMappingFlags Flags =
+ getMapTypeBits(MapType, MapModifiers, IsImplicit,
+ !IsExpressionFirstInfo || RequiresReference,
+ IsCaptureFirstInfo && !RequiresReference,
+ /*IsNonContiguous=*/IsNonContiguous);
if (!IsExpressionFirstInfo) {
// If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
@@ -8359,6 +8387,154 @@
}
}
+ /// Generate the base pointers, section pointers, sizes , map type bits,
+ /// dimension size, offset, count, and strides for the provided map type, map
+ /// modifier, and expression components. \a IsFirstComponent should be set to
+ /// true if the provided set of components is the first associated with a
+ /// capture.
+ void generateInfoForTargetDataComponentList(
+ OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
+ MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
+ MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, MapDimArrayTy &Dims,
+ MapNonContiguousArrayTy &Offsets, MapNonContiguousArrayTy &Counts,
+ MapNonContiguousArrayTy &Strides, StructRangeInfoTy &PartialStruct,
+ bool IsFirstComponentList, bool IsImplicit,
+ ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
+ OverlappedElements = llvm::None) const {
+
+ generateInfoForComponentList(MapType, MapModifiers, Components,
+ BasePointers, Pointers, Sizes, Types, Dims,
+ PartialStruct, IsFirstComponentList,
+ IsImplicit, OverlappedElements, true);
+
+ const ASTContext &Context = CGF.getContext();
+
+ MapValuesArrayTy CurOffsets;
+ MapValuesArrayTy CurCounts;
+ MapValuesArrayTy CurStrides;
+ llvm::Value *CurStride = nullptr;
+
+ // 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.
+ SmallVector<llvm::Value *, 4> DimSizes;
+ for (const auto &Component : Components) {
+ const Expr *AssocExpr = Component.getAssociatedExpression();
+ const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
+ const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+ if (AE || OASE) {
+ QualType Ty;
+ if (OASE)
+ Ty = OMPArraySectionExpr::getBaseOriginalType(OASE->getBase());
+ else
+ Ty = AE->getType();
+ auto *CAT = Context.getAsConstantArrayType(Ty);
+ auto *VAT = Context.getAsVariableArrayType(Ty);
+ // Get element size if we CurStrides is empty.
+ if (CurStrides.empty()) {
+ const Type *ElementType = nullptr;
+ uint64_t ElementTypeSize;
+ if (CAT) {
+ ElementType = CAT->getElementType().getTypePtr();
+ ElementTypeSize =
+ Context.getTypeSizeInChars(ElementType).getQuantity();
+ } else {
+ assert(VAT && "Should be either ConstantArray or VariableArray");
+ ElementType = VAT->getElementType().getTypePtr();
+ ElementTypeSize =
+ Context.getTypeSizeInChars(ElementType).getQuantity();
+ }
+ CurStrides.push_back(
+ llvm::ConstantInt::get(CGF.Int64Ty, ElementTypeSize));
+ }
+ // Get dimension value.
+ llvm::Value *SizeV = nullptr;
+ if (CAT) {
+ llvm::APInt Size = CAT->getSize();
+ SizeV = llvm::ConstantInt::get(CGF.SizeTy, Size);
+ } else {
+ assert(VAT && "Should be either ConstantArray or VariableArray");
+ const Expr *Size = VAT->getSizeExpr();
+ SizeV = CGF.EmitScalarExpr(Size);
+ }
+ SizeV = CGF.Builder.CreateIntCast(SizeV, CGF.Int64Ty,
+ /*IsSigned=*/false);
+ DimSizes.push_back(SizeV);
+ }
+ }
+
+ // Scan the components from the base to the complete expression.
+ auto CI = Components.begin();
+ auto CE = Components.end();
+ auto I = CI;
+ 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 (; I != CE; ++I) {
+ const Expr *AssocExpr = I->getAssociatedExpression();
+ const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
+ const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+
+ if (OASE || AE) {
+ // Offset
+ const Expr *OffsetExpr = nullptr;
+ if (OASE)
+ OffsetExpr = OASE->getLowerBound();
+ else
+ OffsetExpr = AE->getIdx();
+ 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 = nullptr;
+ if (OASE)
+ 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.begin()) {
+ CurStride =
+ CGF.Builder.CreateNUWMul(CurStrides.back(), *std::prev(DI, 1));
+ CurStrides.push_back(CurStride);
+ }
+
+ DI++;
+ }
+ }
+
+ Offsets.push_back(CurOffsets);
+ Counts.push_back(CurCounts);
+ Strides.push_back(CurStrides);
+ }
+
/// Return the adjusted map modifiers if the declaration a capture refers to
/// appears in a first-private clause. This is expected to be used only with
/// directives that start with 'target'.
@@ -8524,7 +8700,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.
@@ -8537,11 +8716,12 @@
OMPClauseMappableExprCommon::MappableExprComponentListRef L,
OpenMPMapClauseKind MapType,
ArrayRef<OpenMPMapModifierKind> MapModifiers,
- bool ReturnDevicePointer, bool IsImplicit) {
+ bool ReturnDevicePointer, bool IsImplicit,
+ bool IsNonContiguous) {
const ValueDecl *VD =
D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
- IsImplicit);
+ IsImplicit, IsNonContiguous);
};
assert(CurDir.is<const OMPExecutableDirective *>() &&
@@ -8550,18 +8730,27 @@
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>())
for (const auto L : C->component_lists()) {
InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifiers(),
- /*ReturnDevicePointer=*/false, C->isImplicit());
+ /*ReturnDevicePointer=*/false, C->isImplicit(),
+ /*IsNonContiguous=*/false);
}
- for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>())
- for (const auto L : C->component_lists()) {
- InfoGen(L.first, L.second, OMPC_MAP_to, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit());
+ for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>()) {
+ auto CI = C->component_lists_begin();
+ auto CE = C->component_lists_end();
+ auto NI = C->non_contiguous_list_begin();
+ for (; CI != CE; ++CI, ++NI) {
+ InfoGen((*CI).first, (*CI).second, OMPC_MAP_to, llvm::None,
+ /*ReturnDevicePointer=*/false, C->isImplicit(), *NI);
}
- for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>())
- for (const auto L : C->component_lists()) {
- InfoGen(L.first, L.second, OMPC_MAP_from, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit());
+ }
+ for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>()) {
+ auto CI = C->component_lists_begin();
+ auto CE = C->component_lists_end();
+ auto NI = C->non_contiguous_list_begin();
+ for (; CI != CE; ++CI, ++NI) {
+ InfoGen((*CI).first, (*CI).second, OMPC_MAP_from, llvm::None,
+ /*ReturnDevicePointer=*/false, C->isImplicit(), *NI);
}
+ }
// Look at the use_device_ptr clause information and mark the existing map
// entries as such. If there is no map information for an entry in the
@@ -8588,7 +8777,8 @@
// Look for the first set of components that refer to it.
if (It != Info.end()) {
auto CI = std::find_if(
- It->second.begin(), It->second.end(), [VD](const MapInfo &MI) {
+ It->second.begin(), It->second.end(),
+ [VD](const MapInfo &MI) {
return MI.Components.back().getAssociatedDeclaration() == VD;
});
// If we found a map entry, signal that the pointer has to be returned
@@ -8611,7 +8801,8 @@
// the pointer into account for the calculation of the range of the
// partial struct.
InfoGen(nullptr, L.second, OMPC_MAP_unknown, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit());
+ /*ReturnDevicePointer=*/false, C->isImplicit(),
+ /*IsNonContiguous=*/false);
DeferredInfo[nullptr].emplace_back(IE, VD);
} else {
llvm::Value *Ptr =
@@ -8634,6 +8825,10 @@
MapValuesArrayTy CurPointers;
MapValuesArrayTy CurSizes;
MapFlagsArrayTy CurTypes;
+ MapDimArrayTy CurDims;
+ MapNonContiguousArrayTy CurOffsets;
+ MapNonContiguousArrayTy CurCounts;
+ MapNonContiguousArrayTy CurStrides;
StructRangeInfoTy PartialStruct;
for (const MapInfo &L : M.second) {
@@ -8642,10 +8837,18 @@
// 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);
+ if (L.IsNonContiguous) {
+ generateInfoForTargetDataComponentList(
+ L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+ CurPointers, CurSizes, CurTypes, CurDims, CurOffsets, CurCounts,
+ CurStrides, PartialStruct, IsFirstComponentList, L.IsImplicit);
+ } else {
+ // Indicate that we do not do the special non-contiguous codegen
+ generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
+ CurBasePointers, CurPointers, CurSizes,
+ CurTypes, CurDims, PartialStruct,
+ IsFirstComponentList, L.IsImplicit);
+ }
// If this entry relates with a device pointer, set the relevant
// declaration and add the 'return pointer' flag.
@@ -8685,15 +8888,24 @@
// 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 `emitTargetDataOffloadingArrays` use it
+ // to skip this one.
+ CurDims.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());
Pointers.append(CurPointers.begin(), CurPointers.end());
Sizes.append(CurSizes.begin(), CurSizes.end());
Types.append(CurTypes.begin(), CurTypes.end());
+ Dims.append(CurDims.begin(), CurDims.end());
+ Offsets.append(CurOffsets.begin(), CurOffsets.end());
+ Counts.append(CurCounts.begin(), CurCounts.end());
+ Strides.append(CurStrides.begin(), CurStrides.end());
}
}
@@ -8722,7 +8934,7 @@
const ValueDecl *VD =
D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
- IsImplicit);
+ IsImplicit, /*IsNonContiguous=*/false);
};
for (const auto *C : CurMapperDir->clauselists()) {
@@ -8743,6 +8955,7 @@
MapValuesArrayTy CurPointers;
MapValuesArrayTy CurSizes;
MapFlagsArrayTy CurTypes;
+ MapDimArrayTy CurDims;
StructRangeInfoTy PartialStruct;
for (const MapInfo &L : M.second) {
@@ -8750,7 +8963,7 @@
"Not expecting declaration with no component lists.");
generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
CurBasePointers, CurPointers, CurSizes,
- CurTypes, PartialStruct,
+ CurTypes, CurDims, PartialStruct,
IsFirstComponentList, L.IsImplicit);
IsFirstComponentList = false;
}
@@ -8869,6 +9082,7 @@
MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
+ MapDimArrayTy &Dims,
StructRangeInfoTy &PartialStruct) const {
assert(!Cap->capturesVariableArrayType() &&
"Not expecting to generate map info for a variable array type!");
@@ -9018,7 +9232,7 @@
OverlappedComponents = Pair.getSecond();
bool IsFirstComponentList = true;
generateInfoForComponentList(MapType, MapModifiers, Components,
- BasePointers, Pointers, Sizes, Types,
+ BasePointers, Pointers, Sizes, Types, Dims,
PartialStruct, IsFirstComponentList,
IsImplicit, OverlappedComponents);
}
@@ -9032,10 +9246,9 @@
std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
auto It = OverlappedData.find(&L);
if (It == OverlappedData.end())
- generateInfoForComponentList(MapType, MapModifiers, Components,
- BasePointers, Pointers, Sizes, Types,
- PartialStruct, IsFirstComponentList,
- IsImplicit);
+ generateInfoForComponentList(
+ MapType, MapModifiers, Components, BasePointers, Pointers, Sizes,
+ Types, Dims, PartialStruct, IsFirstComponentList, IsImplicit);
IsFirstComponentList = false;
}
}
@@ -9045,7 +9258,8 @@
void generateInfoForDeclareTargetLink(MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes,
- MapFlagsArrayTy &Types) const {
+ MapFlagsArrayTy &Types,
+ MapDimArrayTy &Dims) const {
assert(CurDir.is<const OMPExecutableDirective *>() &&
"Expect a executable directive");
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
@@ -9066,7 +9280,7 @@
StructRangeInfoTy PartialStruct;
generateInfoForComponentList(
C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers,
- Pointers, Sizes, Types, PartialStruct,
+ Pointers, Sizes, Types, Dims, PartialStruct,
/*IsFirstComponentList=*/true, C->isImplicit());
assert(!PartialStruct.Base.isValid() &&
"No partial structs for declare target link expected.");
@@ -9160,16 +9374,15 @@
};
} // anonymous namespace
-/// 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.
static void
emitOffloadingArrays(CodeGenFunction &CGF,
MappableExprsHandler::MapBaseValuesArrayTy &BasePointers,
MappableExprsHandler::MapValuesArrayTy &Pointers,
MappableExprsHandler::MapValuesArrayTy &Sizes,
MappableExprsHandler::MapFlagsArrayTy &MapTypes,
- CGOpenMPRuntime::TargetDataInfo &Info) {
+ MappableExprsHandler::MapDimArrayTy &Dims,
+ CGOpenMPRuntime::TargetDataInfo &Info,
+ bool IsNonContiguous = false) {
CodeGenModule &CGM = CGF.CGM;
ASTContext &Ctx = CGF.getContext();
@@ -9212,8 +9425,14 @@
// 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, Dims[I]));
+ } else {
+ ConstSizes.push_back(cast<llvm::Constant>(Sizes[I]));
+ }
+ }
auto *SizesArrayInit = llvm::ConstantArray::get(
llvm::ArrayType::get(CGM.Int64Ty, ConstSizes.size()), ConstSizes);
@@ -9279,6 +9498,87 @@
}
}
+/// 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.
+static void
+emitTargetDataOffloadingArrays(CodeGenFunction &CGF,
+ MappableExprsHandler::MapBaseValuesArrayTy &BasePointers,
+ MappableExprsHandler::MapValuesArrayTy &Pointers,
+ MappableExprsHandler::MapValuesArrayTy &Sizes,
+ MappableExprsHandler::MapFlagsArrayTy &MapTypes,
+ MappableExprsHandler::MapDimArrayTy &Dims,
+ MappableExprsHandler::MapNonContiguousArrayTy &Offsets,
+ MappableExprsHandler::MapNonContiguousArrayTy &Counts,
+ MappableExprsHandler::MapNonContiguousArrayTy &Strides,
+ CGOpenMPRuntime::TargetDataInfo &Info) {
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Dims, Info,
+ true);
+
+ if (Offsets.empty()) return;
+
+ ASTContext &C = CGF.getContext();
+ CodeGenModule &CGM = CGF.CGM;
+
+ // Build an array of struct descriptor_dim and then assign it to offload_args.
+ if (Info.NumberOfPtrs) {
+ // Build struct descriptor_dim {
+ // int64_t offset;
+ // int64_t count;
+ // int64_t stride
+ // };
+ 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.NumberOfPtrs; I < E; ++I) {
+ if (Dims[I] == 0)
+ continue;
+ llvm::APInt Size(/*numBits=*/32, Dims[I]);
+ QualType ArrayTy =
+ C.getConstantArrayType(DimTy, Size, nullptr, ArrayType::Normal, 0);
+ Address DimsAddr = CGF.CreateMemTemp(ArrayTy, "dims");
+ for (unsigned II = 0, EE = 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(Offsets[L][RevIdx], OffsetLVal);
+ // Count
+ LValue CountLVal = CGF.EmitLValueForField(
+ DimsLVal, *std::next(RD->field_begin(), CountFD));
+ CGF.EmitStoreOfScalar(Counts[L][RevIdx], CountLVal);
+ // Stride
+ LValue StrideLVal = CGF.EmitLValueForField(
+ DimsLVal, *std::next(RD->field_begin(), StrideFD));
+ CGF.EmitStoreOfScalar(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, C.getTypeAlignInChars(C.VoidPtrTy));
+ CGF.Builder.CreateStore(DAddr.getPointer(), PAddr);
+ ++L;
+ }
+ }
+}
+
/// Emit the arguments to be passed to the runtime library based on the
/// arrays of pointers, sizes and map types.
static void emitOffloadingArraysArgument(
@@ -9952,6 +10252,7 @@
MappableExprsHandler::MapValuesArrayTy Pointers;
MappableExprsHandler::MapValuesArrayTy Sizes;
MappableExprsHandler::MapFlagsArrayTy MapTypes;
+ MappableExprsHandler::MapDimArrayTy Dims;
// Get mappable expression information.
MappableExprsHandler MEHandler(D, CGF);
@@ -9966,6 +10267,7 @@
MappableExprsHandler::MapValuesArrayTy CurPointers;
MappableExprsHandler::MapValuesArrayTy CurSizes;
MappableExprsHandler::MapFlagsArrayTy CurMapTypes;
+ MappableExprsHandler::MapDimArrayTy CurDims;
MappableExprsHandler::StructRangeInfoTy PartialStruct;
// VLA sizes are passed to the outlined region by copy and do not have map
@@ -9983,7 +10285,8 @@
// If we have any information in the map clause, we use it, otherwise we
// just do a default mapping.
MEHandler.generateInfoForCapture(CI, *CV, CurBasePointers, CurPointers,
- CurSizes, CurMapTypes, PartialStruct);
+ CurSizes, CurMapTypes, CurDims,
+ PartialStruct);
if (CurBasePointers.empty())
MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers,
CurPointers, CurSizes, CurMapTypes);
@@ -10020,11 +10323,12 @@
// Map other list items in the map clause which are not captured variables
// but "declare target link" global variables.
MEHandler.generateInfoForDeclareTargetLink(BasePointers, Pointers, Sizes,
- MapTypes);
+ MapTypes, Dims);
TargetDataInfo Info;
// Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Dims,
+ Info);
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
Info.PointersArray, Info.SizesArray,
Info.MapTypesArray, Info);
@@ -10621,13 +10925,19 @@
MappableExprsHandler::MapValuesArrayTy Pointers;
MappableExprsHandler::MapValuesArrayTy Sizes;
MappableExprsHandler::MapFlagsArrayTy MapTypes;
+ MappableExprsHandler::MapDimArrayTy Dims;
+ MappableExprsHandler::MapNonContiguousArrayTy Offsets;
+ MappableExprsHandler::MapNonContiguousArrayTy Counts;
+ MappableExprsHandler::MapNonContiguousArrayTy Strides;
// Get map clause information.
MappableExprsHandler MCHandler(D, CGF);
- MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+ MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Dims,
+ Offsets, Counts, Strides);
// Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+ emitTargetDataOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes,
+ Dims, Offsets, Counts, Strides, Info);
llvm::Value *BasePointersArrayArg = nullptr;
llvm::Value *PointersArrayArg = nullptr;
@@ -10857,14 +11167,20 @@
MappableExprsHandler::MapValuesArrayTy Pointers;
MappableExprsHandler::MapValuesArrayTy Sizes;
MappableExprsHandler::MapFlagsArrayTy MapTypes;
+ MappableExprsHandler::MapDimArrayTy Dims;
+ MappableExprsHandler::MapNonContiguousArrayTy Offsets;
+ MappableExprsHandler::MapNonContiguousArrayTy Counts;
+ MappableExprsHandler::MapNonContiguousArrayTy Strides;
// Get map clause information.
MappableExprsHandler MEHandler(D, CGF);
- MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+ MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Dims,
+ Offsets, Counts, Strides);
TargetDataInfo Info;
// Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+ emitTargetDataOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes,
+ Dims, Offsets, Counts, Strides, Info);
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
Info.PointersArray, Info.SizesArray,
Info.MapTypesArray, Info);
Index: clang/lib/AST/OpenMPClause.cpp
===================================================================
--- clang/lib/AST/OpenMPClause.cpp
+++ clang/lib/AST/OpenMPClause.cpp
@@ -986,7 +986,8 @@
OMPMapClause *OMPMapClause::Create(
const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
- MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+ MappableExprComponentListsRef ComponentLists,
+ ArrayRef<bool> NonContiguousList, ArrayRef<Expr *> UDMapperRefs,
ArrayRef<OpenMPMapModifierKind> MapModifiers,
ArrayRef<SourceLocation> MapModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId,
@@ -1002,15 +1003,17 @@
// user-defined mapper for each clause list entry.
// NumUniqueDeclarations x ValueDecl* - unique base declarations associated
// with each component list.
+ // NumComponentLists x bool - number of non-contiguous attribute.
// (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
// number of lists for each unique declaration and the size of each component
// list.
// NumComponents x MappableComponent - the total of all the components in all
// the lists.
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
OMPMapClause *Clause = new (Mem)
@@ -1019,7 +1022,7 @@
Clause->setVarRefs(Vars);
Clause->setUDMapperRefs(UDMapperRefs);
- Clause->setClauseInfo(Declarations, ComponentLists);
+ Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
Clause->setMapType(Type);
Clause->setMapLoc(TypeLoc);
return Clause;
@@ -1029,9 +1032,9 @@
OMPMapClause::CreateEmpty(const ASTContext &C,
const OMPMappableExprListSizeTy &Sizes) {
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
- 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumVars,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
return new (Mem) OMPMapClause(Sizes);
@@ -1040,7 +1043,8 @@
OMPToClause *OMPToClause::Create(
const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
- MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+ MappableExprComponentListsRef ComponentLists,
+ ArrayRef<bool> NonContiguousList, ArrayRef<Expr *> UDMapperRefs,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
OMPMappableExprListSizeTy Sizes;
Sizes.NumVars = Vars.size();
@@ -1053,15 +1057,17 @@
// user-defined mapper for each clause list entry.
// NumUniqueDeclarations x ValueDecl* - unique base declarations associated
// with each component list.
+ // NumComponentLists x bool - number of non-contiguous attribute.
// (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
// number of lists for each unique declaration and the size of each component
// list.
// NumComponents x MappableComponent - the total of all the components in all
// the lists.
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
@@ -1069,16 +1075,18 @@
Clause->setVarRefs(Vars);
Clause->setUDMapperRefs(UDMapperRefs);
- Clause->setClauseInfo(Declarations, ComponentLists);
+ Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
+
return Clause;
}
OMPToClause *OMPToClause::CreateEmpty(const ASTContext &C,
const OMPMappableExprListSizeTy &Sizes) {
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
return new (Mem) OMPToClause(Sizes);
@@ -1087,7 +1095,9 @@
OMPFromClause *OMPFromClause::Create(
const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
- MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+ MappableExprComponentListsRef ComponentLists,
+ ArrayRef<bool> NonContiguousList,
+ ArrayRef<Expr *> UDMapperRefs,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
OMPMappableExprListSizeTy Sizes;
Sizes.NumVars = Vars.size();
@@ -1100,15 +1110,17 @@
// user-defined mapper for each clause list entry.
// NumUniqueDeclarations x ValueDecl* - unique base declarations associated
// with each component list.
+ // NumComponentLists x bool - number of non-contiguous attribute.
// (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
// number of lists for each unique declaration and the size of each component
// list.
// NumComponents x MappableComponent - the total of all the components in all
// the lists.
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
@@ -1117,7 +1129,8 @@
Clause->setVarRefs(Vars);
Clause->setUDMapperRefs(UDMapperRefs);
- Clause->setClauseInfo(Declarations, ComponentLists);
+ Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
+
return Clause;
}
@@ -1125,9 +1138,10 @@
OMPFromClause::CreateEmpty(const ASTContext &C,
const OMPMappableExprListSizeTy &Sizes) {
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
return new (Mem) OMPFromClause(Sizes);
@@ -1161,15 +1175,17 @@
// list entry and an equal number of private copies and inits.
// NumUniqueDeclarations x ValueDecl* - unique base declarations associated
// with each component list.
+ // NumComponentLists x bool - number of non-contiguous attribute.
// (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
// number of lists for each unique declaration and the size of each component
// list.
// NumComponents x MappableComponent - the total of all the components in all
// the lists.
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
3 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
@@ -1178,7 +1194,8 @@
Clause->setVarRefs(Vars);
Clause->setPrivateCopies(PrivateVars);
Clause->setInits(Inits);
- Clause->setClauseInfo(Declarations, ComponentLists);
+ SmallVector<bool, 4> NonContiguousList(Declarations.size(), false);
+ Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
return Clause;
}
@@ -1186,9 +1203,10 @@
OMPUseDevicePtrClause::CreateEmpty(const ASTContext &C,
const OMPMappableExprListSizeTy &Sizes) {
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
3 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
return new (Mem) OMPUseDevicePtrClause(Sizes);
@@ -1210,22 +1228,24 @@
// entry.
// NumUniqueDeclarations x ValueDecl* - unique base declarations associated
// with each component list.
+ // NumComponentLists x bool - number of non-contiguous attribute.
// (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
// number of lists for each unique declaration and the size of each component
// list.
// NumComponents x MappableComponent - the total of all the components in all
// the lists.
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
- Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
OMPIsDevicePtrClause *Clause = new (Mem) OMPIsDevicePtrClause(Locs, Sizes);
Clause->setVarRefs(Vars);
- Clause->setClauseInfo(Declarations, ComponentLists);
+ SmallVector<bool, 4> NonContiguousList(Declarations.size(), false);
+ Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
return Clause;
}
@@ -1233,9 +1253,9 @@
OMPIsDevicePtrClause::CreateEmpty(const ASTContext &C,
const OMPMappableExprListSizeTy &Sizes) {
void *Mem = C.Allocate(
- totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
- Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumComponentLists,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
return new (Mem) OMPIsDevicePtrClause(Sizes);
Index: clang/include/clang/AST/OpenMPClause.h
===================================================================
--- clang/include/clang/AST/OpenMPClause.h
+++ clang/include/clang/AST/OpenMPClause.h
@@ -4886,10 +4886,35 @@
std::copy(Components.begin(), Components.end(), getComponentsRef().begin());
}
+ /// Get the non-contiguous attribute per declaration that are in the trailing
+ /// objects of the class.
+ MutableArrayRef<bool> getNonContiguousListsRef() {
+ return MutableArrayRef<bool>(
+ static_cast<T *>(this)->template getTrailingObjects<bool>(),
+ NumComponentLists);
+ }
+
+ /// Get the non-contiguous attribute per declaration that are in the trailing
+ /// objects of the class.
+ ArrayRef<bool> getNonContiguousListsRef() const {
+ return ArrayRef<bool>(
+ static_cast<const T *>(this)->template getTrailingObjects<bool>(),
+ NumComponentLists);
+ }
+
+ /// Set the non-contiguous attribute per declaration that are in the trailing
+ /// objects of the class.
+ void setNonContiguousLists(ArrayRef<bool> NLs) {
+ assert(NLs.size() == NumComponentLists &&
+ "Unexpected amount of list numbers.");
+ std::copy(NLs.begin(), NLs.end(), getNonContiguousListsRef().begin());
+ }
+
/// Fill the clause information from the list of declarations and
/// associated component lists.
void setClauseInfo(ArrayRef<ValueDecl *> Declarations,
- MappableExprComponentListsRef ComponentLists) {
+ MappableExprComponentListsRef ComponentLists,
+ ArrayRef<bool> NonContiguousList) {
// Perform some checks to make sure the data sizes are consistent with the
// information available when the clause was created.
assert(getUniqueDeclarationsTotalNumber(Declarations) ==
@@ -4901,6 +4926,8 @@
"Declaration and component lists size is not consistent!");
assert(Declarations.size() == NumComponentLists &&
"Unexpected declaration and component lists size!");
+ assert(NonContiguousList.size() == ComponentLists.size() &&
+ "Unexpected NonContiguousList size");
// Organize the components by declaration and retrieve the original
// expression. Original expressions are always the first component of the
@@ -4960,6 +4987,9 @@
CI = std::copy(C.begin(), C.end(), CI);
}
}
+
+ std::copy(NonContiguousList.begin(), NonContiguousList.end(),
+ getNonContiguousListsRef().begin());
}
/// Set the nested name specifier of associated user-defined mapper.
@@ -5221,6 +5251,34 @@
return const_all_components_range(A.begin(), A.end());
}
+ using non_contiguous_list_iterator = MutableArrayRef<bool>::iterator;
+ using non_contiguous_list_const_iterator = ArrayRef<bool>::iterator;
+ using non_contiguous_list_range =
+ llvm::iterator_range<non_contiguous_list_iterator>;
+ using non_contiguous_list_const_range =
+ llvm::iterator_range<non_contiguous_list_const_iterator>;
+
+ non_contiguous_list_iterator non_contiguous_list_begin() {
+ return getNonContiguousListsRef().begin();
+ }
+ non_contiguous_list_iterator non_contiguous_list_end() {
+ return getNonContiguousListsRef().end();
+ }
+ non_contiguous_list_const_iterator non_contiguous_list_begin() const {
+ return getNonContiguousListsRef().begin();
+ }
+ non_contiguous_list_const_iterator non_contiguous_list_end() const {
+ return getNonContiguousListsRef().end();
+ }
+ non_contiguous_list_range non_contiguous_lists() {
+ return non_contiguous_list_range(non_contiguous_list_begin(),
+ non_contiguous_list_end());
+ }
+ non_contiguous_list_const_range non_contiguous_lists() const {
+ return non_contiguous_list_const_range(non_contiguous_list_begin(),
+ non_contiguous_list_end());
+ }
+
using mapperlist_iterator = MutableArrayRef<Expr *>::iterator;
using mapperlist_const_iterator = ArrayRef<const Expr *>::iterator;
using mapperlist_range = llvm::iterator_range<mapperlist_iterator>;
@@ -5251,10 +5309,11 @@
/// \endcode
/// In this example directive '#pragma omp target' has clause 'map'
/// with the variables 'a' and 'b'.
-class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
- private llvm::TrailingObjects<
- OMPMapClause, Expr *, ValueDecl *, unsigned,
- OMPClauseMappableExprCommon::MappableComponent> {
+class OMPMapClause final
+ : public OMPMappableExprListClause<OMPMapClause>,
+ private llvm::TrailingObjects<
+ OMPMapClause, Expr *, ValueDecl *, bool, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent> {
friend class OMPClauseReader;
friend OMPMappableExprListClause;
friend OMPVarListClause;
@@ -5270,6 +5329,9 @@
size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
return getUniqueDeclarationsNum();
}
+ size_t numTrailingObjects(OverloadToken<bool>) const {
+ return getTotalComponentListNum();
+ }
size_t numTrailingObjects(OverloadToken<unsigned>) const {
return getUniqueDeclarationsNum() + getTotalComponentListNum();
}
@@ -5403,7 +5465,7 @@
Create(const ASTContext &C, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists,
- ArrayRef<Expr *> UDMapperRefs,
+ ArrayRef<bool> NonContiguousList, ArrayRef<Expr *> UDMapperRefs,
ArrayRef<OpenMPMapModifierKind> MapModifiers,
ArrayRef<SourceLocation> MapModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId,
@@ -6206,7 +6268,7 @@
/// with the variables 'a' and 'b'.
class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
private llvm::TrailingObjects<
- OMPToClause, Expr *, ValueDecl *, unsigned,
+ OMPToClause, Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent> {
friend class OMPClauseReader;
friend OMPMappableExprListClause;
@@ -6254,6 +6316,9 @@
size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
return getUniqueDeclarationsNum();
}
+ size_t numTrailingObjects(OverloadToken<bool>) const {
+ return getTotalComponentListNum();
+ }
size_t numTrailingObjects(OverloadToken<unsigned>) const {
return getUniqueDeclarationsNum() + getTotalComponentListNum();
}
@@ -6277,6 +6342,7 @@
ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists,
+ ArrayRef<bool> IsNonContiguousList,
ArrayRef<Expr *> UDMapperRefs,
NestedNameSpecifierLoc UDMQualifierLoc,
DeclarationNameInfo MapperId);
@@ -6325,7 +6391,7 @@
class OMPFromClause final
: public OMPMappableExprListClause<OMPFromClause>,
private llvm::TrailingObjects<
- OMPFromClause, Expr *, ValueDecl *, unsigned,
+ OMPFromClause, Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent> {
friend class OMPClauseReader;
friend OMPMappableExprListClause;
@@ -6373,6 +6439,9 @@
size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
return getUniqueDeclarationsNum();
}
+ size_t numTrailingObjects(OverloadToken<bool>) const {
+ return getTotalComponentListNum();
+ }
size_t numTrailingObjects(OverloadToken<unsigned>) const {
return getUniqueDeclarationsNum() + getTotalComponentListNum();
}
@@ -6396,6 +6465,7 @@
ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists,
+ ArrayRef<bool> NonContiguousList,
ArrayRef<Expr *> UDMapperRefs,
NestedNameSpecifierLoc UDMQualifierLoc,
DeclarationNameInfo MapperId);
@@ -6444,7 +6514,7 @@
class OMPUseDevicePtrClause final
: public OMPMappableExprListClause<OMPUseDevicePtrClause>,
private llvm::TrailingObjects<
- OMPUseDevicePtrClause, Expr *, ValueDecl *, unsigned,
+ OMPUseDevicePtrClause, Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent> {
friend class OMPClauseReader;
friend OMPMappableExprListClause;
@@ -6485,6 +6555,9 @@
size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
return getUniqueDeclarationsNum();
}
+ size_t numTrailingObjects(OverloadToken<bool>) const {
+ return getTotalComponentListNum();
+ }
size_t numTrailingObjects(OverloadToken<unsigned>) const {
return getUniqueDeclarationsNum() + getTotalComponentListNum();
}
@@ -6608,7 +6681,7 @@
class OMPIsDevicePtrClause final
: public OMPMappableExprListClause<OMPIsDevicePtrClause>,
private llvm::TrailingObjects<
- OMPIsDevicePtrClause, Expr *, ValueDecl *, unsigned,
+ OMPIsDevicePtrClause, Expr *, ValueDecl *, bool, unsigned,
OMPClauseMappableExprCommon::MappableComponent> {
friend class OMPClauseReader;
friend OMPMappableExprListClause;
@@ -6648,6 +6721,9 @@
size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
return getUniqueDeclarationsNum();
}
+ size_t numTrailingObjects(OverloadToken<bool>) const {
+ return getTotalComponentListNum();
+ }
size_t numTrailingObjects(OverloadToken<unsigned>) const {
return getUniqueDeclarationsNum() + getTotalComponentListNum();
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
