lildmh updated this revision to Diff 198050.
lildmh added a comment.
Fix code format
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D59474/new/
https://reviews.llvm.org/D59474
Files:
include/clang/AST/GlobalDecl.h
lib/AST/ASTContext.cpp
lib/CodeGen/CGDecl.cpp
lib/CodeGen/CGOpenMPRuntime.cpp
lib/CodeGen/CGOpenMPRuntime.h
lib/CodeGen/CodeGenModule.h
lib/CodeGen/ModuleBuilder.cpp
test/OpenMP/declare_mapper_codegen.cpp
Index: test/OpenMP/declare_mapper_codegen.cpp
===================================================================
--- test/OpenMP/declare_mapper_codegen.cpp
+++ test/OpenMP/declare_mapper_codegen.cpp
@@ -1,92 +1,770 @@
-///==========================================================================///
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s
-// RUN: %clang_cc1 -fopenmp -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-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s
-// RUN: %clang_cc1 -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -fopenmp-simd -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-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -fopenmp-simd -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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s
-
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
+///==========================================================================///
+// RUN: %clang_cc1 -DCK0 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK0 --check-prefix CK0-64 %s
+// RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK0 --check-prefix CK0-64 %s
+// RUN: %clang_cc1 -DCK0 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK0 --check-prefix CK0-32 %s
+// RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK0 --check-prefix CK0-32 %s
+
+// RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+
+#ifdef CK0
+
+// CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
+// CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i64 16]
+// CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i32 8]
+// CK0: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
+// CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 16]
+// CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 8]
+// CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 16]
+// CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 8]
+// CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+
class C {
public:
int a;
+ double *b;
};
-#pragma omp declare mapper(id: C s) map(s.a)
+#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
+
+// Synchronous version of mapper function.
+// CK0-LABEL: define {{.*}}i32 @.omp_mapper.class_C.id{{.*}}(i64, i8*, i8*, i{{64|32}}, i64)
+// CK0-DAG: store i64 %0, i64* [[DIDADDR:%[^,]+]]
+// CK0-DAG: store i[[sz:64|32]] %3, i{{64|32}}* [[SIZEADDR:%[^,]+]]
+// CK0-DAG: store i64 %4, i64* [[TYPEADDR:%[^,]+]]
+// CK0-DAG: store i8* %1, i8** [[BPTRADDR:%[^,]+]]
+// CK0-DAG: store i8* %2, i8** [[VPTRADDR:%[^,]+]]
+// CK0-DAG: store i32 0, i32* %retval
+// CK0-DAG: [[SIZE:%.+]] = load i[[sz]], i[[sz]]* [[SIZEADDR]]
+// CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
+// CK0-DAG: [[DID:%.+]] = load i64, i64* [[DIDADDR]]
+// CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
+// CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i[[sz]] [[SIZE]]
+// CK0: [[ISARRAY:%.+]] = icmp sge i[[sz]] [[SIZE]], 1
+// CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
+
+// CK0: [[INITEVALDEL]]
+// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK0: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
+// CK0: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK0: [[INIT]]
+// CK0-64: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 16
+// CK0-32: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 8
+// CK0-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[ISIZEADDR:[^,]+]]
+// CK0-DAG: [[ISIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[ISIZE:%.+]], i32 0, i32 0
+// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: store i64 [[ITYPE]], i64* [[ITYPEADDR:[^,]+]]
+// CK0-DAG: [[ITYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[ITYPE:%.+]], i32 0, i32 0
+// CK0: [[IRES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[ISIZEADDR]], i64* [[ITYPEADDR]], i8** null)
+// CK0: [[ISINITERR:%.+]] = icmp ne i32 [[IRES]], 0
+// CK0: br i1 [[ISINITERR]], label %[[INITERR:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK0: [[INITERR]]
+// CK0: store i32 [[IRES]], i32* %retval
+// CK0: br label %[[DONE:[^,]+]]
+
+// CK0: [[LHEAD]]
+// CK0: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
+// CK0: br i1 [[ISEMPTY]], label %[[DONE]], label %[[LBODY:[^,]+]]
+// CK0: [[LBODY]]
+// CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK0: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
+// CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
+// CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
+// CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
+// CK0-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]]
+// CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz]] 0
+// CK0-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1
+// CK0-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
+// CK0-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8*
+// CK0-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
+// CK0-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
+// CK0-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
+// CK0-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK0-32-DAG: [[CUSIZE32:%.+]] = trunc i64 [[CUSIZE]] to i32
+// CK0-64-DAG: store i[[sz]] [[CUSIZE]], i[[sz]]* [[SIZEADDR0:%[^,]+]]
+// CK0-32-DAG: store i[[sz]] [[CUSIZE32]], i[[sz]]* [[SIZEADDR0:%[^,]+]]
+// CK0-DAG: [[SIZEADDR0]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR:%[^,]+]], i32 0, i32 0
+// CK0-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR0BC:%[^,]+]]
+// CK0-DAG: [[BPTRADDR0BC]] = bitcast i8** [[BPTRADDR0:%.+]] to %class.C**
+// CK0-DAG: [[BPTRADDR0]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR:%[^,]+]], i32 0, i32 0
+// CK0-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR0BC:%[^,]+]]
+// CK0-DAG: [[PTRADDR0BC]] = bitcast i8** [[PTRADDR0:%.+]] to i32**
+// CK0-DAG: [[PTRADDR0]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR:%[^,]+]], i32 0, i32 0
+// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK0-DAG: [[ALLOC]]
+// CK0-DAG: br label %[[TYEND:[^,]+]]
+// CK0-DAG: [[ALLOCELSE]]
+// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK0-DAG: [[TO]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TOELSE]]
+// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK0-DAG: [[FROM]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TYEND]]
+// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ 32, %[[ALLOC]] ], [ 32, %[[TO]] ], [ 32, %[[FROM]] ], [ 32, %[[TOELSE]] ]
+// CK0-DAG: store i64 [[PHITYPE0]], i64* [[TYPEADDR0:%[^,]+]]
+// CK0-DAG: [[TYPEADDR0]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR:%[^,]+]], i32 0, i32 0
+// CK0-DAG: [[BPTRADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 1
+// CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast i8** [[BPTRADDR1]] to %class.C**
+// CK0-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR1BC]]
+// CK0-DAG: [[PTRADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 1
+// CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i8** [[PTRADDR1]] to i32**
+// CK0-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR1BC]]
+// CK0-DAG: [[SIZEADDR1:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 1
+// CK0-DAG: store i[[sz]] 4, i[[sz]]* [[SIZEADDR1]]
+// CK0-DAG: [[TYPEADDR1:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 1
+// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK0-DAG: [[ALLOC]]
+// CK0-DAG: br label %[[TYEND:[^,]+]]
+// CK0-DAG: [[ALLOCELSE]]
+// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK0-DAG: [[TO]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TOELSE]]
+// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK0-DAG: [[FROM]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TYEND]]
+// 0x1,000,000,000; 0x1,000,000,001; 0x1,000,000,002; 0x1,000,000,003
+// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ 281474976710656, %[[ALLOC]] ], [ 281474976710657, %[[TO]] ], [ 281474976710658, %[[FROM]] ], [ 281474976710659, %[[TOELSE]] ]
+// CK0-DAG: store i64 [[TYPE1]], i64* [[TYPEADDR1]]
+// CK0-DAG: [[BPTRADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 2
+// CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast i8** [[BPTRADDR2]] to double***
+// CK0-DAG: store double** [[BBEGIN]], double*** [[BPTRADDR2BC]]
+// CK0-DAG: [[PTRADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 2
+// CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast i8** [[PTRADDR2]] to double**
+// CK0-DAG: store double* [[BARRBEGINGEP]], double** [[PTRADDR2BC]]
+// CK0-DAG: [[SIZEADDR2:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 2
+// CK0-DAG: store i[[sz]] 16, i[[sz]]* [[SIZEADDR2]]
+// CK0-DAG: [[TYPEADDR2:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 2
+// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK0-DAG: [[ALLOC]]
+// CK0-DAG: br label %[[TYEND:[^,]+]]
+// CK0-DAG: [[ALLOCELSE]]
+// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK0-DAG: [[TO]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TOELSE]]
+// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK0-DAG: [[FROM]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TYEND]]
+// 0x1,000,000,010; 0x1,000,000,011; 0x1,000,000,012; 0x1,000,000,013
+// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ 281474976710672, %[[ALLOC]] ], [ 281474976710673, %[[TO]] ], [ 281474976710674, %[[FROM]] ], [ 281474976710675, %[[TOELSE]] ]
+// CK0-DAG: store i64 [[TYPE2]], i64* [[TYPEADDR2]]
+// CK0-DAG: [[ARGBPTR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 0
+// CK0-DAG: [[ARGPTR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 0
+// CK0-DAG: [[ARGSIZE:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 0
+// CK0-DAG: [[ARGTYPE:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 0
+// CK0: [[RES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 3, i8** [[ARGBPTR]], i8** [[ARGPTR]], i[[sz]]* [[ARGSIZE]], i64* [[ARGTYPE]], i8** null)
+// CK0: [[ISERR:%.+]] = icmp ne i32 [[RES]], 0
+// CK0: br i1 [[ISERR]], label %[[LERR:[^,]+]], label %[[LCORRECT]]
+// CK0: [[LERR]]
+// CK0: store i32 [[RES]], i32* %retval
+// CK0: br label %[[DONE]]
+// CK0: [[LCORRECT]]
+// CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
+// CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
+// CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
-// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l54.region_id = weak constant i8 0
+// CK0: [[LEXIT]]
+// CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
+// CK0: [[EVALDEL]]
+// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
+// CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
+// CK0: [[DEL]]
+// CK0-64: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 16
+// CK0-32: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 8
+// CK0-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[DSIZEADDR:%[^,]+]]
+// CK0-DAG: [[DSIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[DSIZE:%.+]], i32 0, i32 0
+// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: store i64 [[DTYPE]], i64* [[DTYPEADDR:%[^,]+]]
+// CK0-DAG: [[DTYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[DTYPE:%.+]], i32 0, i32 0
+// CK0: [[DRES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[DSIZEADDR]], i64* [[DTYPEADDR]], i8** null)
+// CK0: [[ISDELERR:%.+]] = icmp ne i32 [[DRES]], 0
+// CK0: br i1 [[ISDELERR]], label %[[DELERR:[^,]+]], label %[[DONE]]
+// CK0: [[DELERR]]
+// CK0: store i32 [[DRES]], i32* %retval
+// CK0: br label %[[DONE]]
+// CK0: [[DONE]]
+// CK0: [[RET:%.+]] = load i32, i32* %retval
+// CK0: ret i32 [[RET]]
-// CHECK: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// CHECK: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
-// CHECK: [[TSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CHECK: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
-// CHECK: [[FSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CHECK: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
-// CHECK-LABEL: foo{{.*}}(
+// Asynchronous version of mapper function.
+// CK0-LABEL: define {{.*}}i32 @.omp_mapper.class_C.id{{.*}}nowait{{.*}}(i64, i8*, i8*, i{{64|32}}, i64)
+// CK0-DAG: store i64 %0, i64* [[DIDADDR:%[^,]+]]
+// CK0-DAG: store i[[sz:64|32]] %3, i{{64|32}}* [[SIZEADDR:%[^,]+]]
+// CK0-DAG: store i64 %4, i64* [[TYPEADDR:%[^,]+]]
+// CK0-DAG: store i8* %1, i8** [[BPTRADDR:%[^,]+]]
+// CK0-DAG: store i8* %2, i8** [[VPTRADDR:%[^,]+]]
+// CK0-DAG: store i32 0, i32* %retval
+// CK0-DAG: [[SIZE:%.+]] = load i[[sz]], i[[sz]]* [[SIZEADDR]]
+// CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
+// CK0-DAG: [[DID:%.+]] = load i64, i64* [[DIDADDR]]
+// CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
+// CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i[[sz]] [[SIZE]]
+// CK0: [[ISARRAY:%.+]] = icmp sge i[[sz]] [[SIZE]], 1
+// CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
+
+// CK0: [[INITEVALDEL]]
+// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK0: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
+// CK0: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK0: [[INIT]]
+// CK0-64: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 16
+// CK0-32: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 8
+// CK0-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[ISIZEADDR:[^,]+]]
+// CK0-DAG: [[ISIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[ISIZE:%.+]], i32 0, i32 0
+// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: store i64 [[ITYPE]], i64* [[ITYPEADDR:[^,]+]]
+// CK0-DAG: [[ITYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[ITYPE:%.+]], i32 0, i32 0
+// CK0: [[IRES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[ISIZEADDR]], i64* [[ITYPEADDR]], i8** null)
+// CK0: [[ISINITERR:%.+]] = icmp ne i32 [[IRES]], 0
+// CK0: br i1 [[ISINITERR]], label %[[INITERR:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK0: [[INITERR]]
+// CK0: store i32 [[IRES]], i32* %retval
+// CK0: br label %[[DONE:[^,]+]]
+
+// CK0: [[LHEAD]]
+// CK0: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
+// CK0: br i1 [[ISEMPTY]], label %[[DONE]], label %[[LBODY:[^,]+]]
+// CK0: [[LBODY]]
+// CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK0: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
+// CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
+// CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
+// CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
+// CK0-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]]
+// CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz]] 0
+// CK0-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1
+// CK0-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
+// CK0-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8*
+// CK0-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
+// CK0-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
+// CK0-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
+// CK0-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK0-32-DAG: [[CUSIZE32:%.+]] = trunc i64 [[CUSIZE]] to i32
+// CK0-64-DAG: store i[[sz]] [[CUSIZE]], i[[sz]]* [[SIZEADDR0:%[^,]+]]
+// CK0-32-DAG: store i[[sz]] [[CUSIZE32]], i[[sz]]* [[SIZEADDR0:%[^,]+]]
+// CK0-DAG: [[SIZEADDR0]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR:%[^,]+]], i32 0, i32 0
+// CK0-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR0BC:%[^,]+]]
+// CK0-DAG: [[BPTRADDR0BC]] = bitcast i8** [[BPTRADDR0:%.+]] to %class.C**
+// CK0-DAG: [[BPTRADDR0]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR:%[^,]+]], i32 0, i32 0
+// CK0-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR0BC:%[^,]+]]
+// CK0-DAG: [[PTRADDR0BC]] = bitcast i8** [[PTRADDR0:%.+]] to i32**
+// CK0-DAG: [[PTRADDR0]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR:%[^,]+]], i32 0, i32 0
+// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK0-DAG: [[ALLOC]]
+// CK0-DAG: br label %[[TYEND:[^,]+]]
+// CK0-DAG: [[ALLOCELSE]]
+// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK0-DAG: [[TO]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TOELSE]]
+// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK0-DAG: [[FROM]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TYEND]]
+// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ 32, %[[ALLOC]] ], [ 32, %[[TO]] ], [ 32, %[[FROM]] ], [ 32, %[[TOELSE]] ]
+// CK0-DAG: store i64 [[PHITYPE0]], i64* [[TYPEADDR0:%[^,]+]]
+// CK0-DAG: [[TYPEADDR0]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR:%[^,]+]], i32 0, i32 0
+// CK0-DAG: [[BPTRADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 1
+// CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast i8** [[BPTRADDR1]] to %class.C**
+// CK0-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR1BC]]
+// CK0-DAG: [[PTRADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 1
+// CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i8** [[PTRADDR1]] to i32**
+// CK0-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR1BC]]
+// CK0-DAG: [[SIZEADDR1:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 1
+// CK0-DAG: store i[[sz]] 4, i[[sz]]* [[SIZEADDR1]]
+// CK0-DAG: [[TYPEADDR1:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 1
+// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK0-DAG: [[ALLOC]]
+// CK0-DAG: br label %[[TYEND:[^,]+]]
+// CK0-DAG: [[ALLOCELSE]]
+// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK0-DAG: [[TO]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TOELSE]]
+// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK0-DAG: [[FROM]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TYEND]]
+// 0x1,000,000,000; 0x1,000,000,001; 0x1,000,000,002; 0x1,000,000,003
+// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ 281474976710656, %[[ALLOC]] ], [ 281474976710657, %[[TO]] ], [ 281474976710658, %[[FROM]] ], [ 281474976710659, %[[TOELSE]] ]
+// CK0-DAG: store i64 [[TYPE1]], i64* [[TYPEADDR1]]
+// CK0-DAG: [[BPTRADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 2
+// CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast i8** [[BPTRADDR2]] to double***
+// CK0-DAG: store double** [[BBEGIN]], double*** [[BPTRADDR2BC]]
+// CK0-DAG: [[PTRADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 2
+// CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast i8** [[PTRADDR2]] to double**
+// CK0-DAG: store double* [[BARRBEGINGEP]], double** [[PTRADDR2BC]]
+// CK0-DAG: [[SIZEADDR2:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 2
+// CK0-DAG: store i[[sz]] 16, i[[sz]]* [[SIZEADDR2]]
+// CK0-DAG: [[TYPEADDR2:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 2
+// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK0-DAG: [[ALLOC]]
+// CK0-DAG: br label %[[TYEND:[^,]+]]
+// CK0-DAG: [[ALLOCELSE]]
+// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK0-DAG: [[TO]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TOELSE]]
+// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK0-DAG: [[FROM]]
+// CK0-DAG: br label %[[TYEND]]
+// CK0-DAG: [[TYEND]]
+// 0x1,000,000,010; 0x1,000,000,011; 0x1,000,000,012; 0x1,000,000,013
+// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ 281474976710672, %[[ALLOC]] ], [ 281474976710673, %[[TO]] ], [ 281474976710674, %[[FROM]] ], [ 281474976710675, %[[TOELSE]] ]
+// CK0-DAG: store i64 [[TYPE2]], i64* [[TYPEADDR2]]
+// CK0-DAG: [[ARGBPTR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 0
+// CK0-DAG: [[ARGPTR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 0
+// CK0-DAG: [[ARGSIZE:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 0
+// CK0-DAG: [[ARGTYPE:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 0
+// CK0: [[RES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 3, i8** [[ARGBPTR]], i8** [[ARGPTR]], i[[sz]]* [[ARGSIZE]], i64* [[ARGTYPE]], i8** null)
+// CK0: [[ISERR:%.+]] = icmp ne i32 [[RES]], 0
+// CK0: br i1 [[ISERR]], label %[[LERR:[^,]+]], label %[[LCORRECT]]
+// CK0: [[LERR]]
+// CK0: store i32 [[RES]], i32* %retval
+// CK0: br label %[[DONE]]
+// CK0: [[LCORRECT]]
+// CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
+// CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
+// CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
+
+// CK0: [[LEXIT]]
+// CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
+// CK0: [[EVALDEL]]
+// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
+// CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
+// CK0: [[DEL]]
+// CK0-64: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 16
+// CK0-32: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 8
+// CK0-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[DSIZEADDR:%[^,]+]]
+// CK0-DAG: [[DSIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[DSIZE:%.+]], i32 0, i32 0
+// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: store i64 [[DTYPE]], i64* [[DTYPEADDR:%[^,]+]]
+// CK0-DAG: [[DTYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[DTYPE:%.+]], i32 0, i32 0
+// CK0: [[DRES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[DSIZEADDR]], i64* [[DTYPEADDR]], i8** null)
+// CK0: [[ISDELERR:%.+]] = icmp ne i32 [[DRES]], 0
+// CK0: br i1 [[ISDELERR]], label %[[DELERR:[^,]+]], label %[[DONE]]
+// CK0: [[DELERR]]
+// CK0: store i32 [[DRES]], i32* %retval
+// CK0: br label %[[DONE]]
+// CK0: [[DONE]]
+// CK0: [[RET:%.+]] = load i32, i32* %retval
+// CK0: ret i32 [[RET]]
+
+
+// CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
void foo(int a){
int i = a;
C c;
c.a = a;
- // CHECK-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
- // CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
- // CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
- // CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
- // CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
- // CHECK-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
- // CHECK-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
- // CHECK-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
- // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
- // CHECK: call void [[KERNEL:@.+]](%class.C* [[VAL]])
+ // CK0-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
+ // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
+ // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
+ // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
+ // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
+ // CK0: call void [[KERNEL:@.+]](%class.C* [[VAL]])
#pragma omp target map(mapper(id),tofrom: c)
{
++c.a;
}
- // CHECK-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}})
- // CHECK-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
- // CHECK-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
- // CHECK-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
- // CHECK-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
- // CHECK-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
- // CHECK-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
- // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]]
- // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
+ // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}})
+ // CK0-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
+ // CK0-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
+ // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]]
+ // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
#pragma omp target update to(mapper(id): c)
- // CHECK-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}})
- // CHECK-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
- // CHECK-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
- // CHECK-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
- // CHECK-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
- // CHECK-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
- // CHECK-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
- // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
- // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
+ // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}})
+ // CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
+ // CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
+ // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
+ // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
#pragma omp target update from(mapper(id): c)
}
-// CHECK: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]])
-// CHECK: [[ADDR:%.+]] = alloca %class.C*,
-// CHECK: store %class.C* [[ARG]], %class.C** [[ADDR]]
-// CHECK: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
-// CHECK: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
-// CHECK: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
-// CHECK: {{.+}} = add nsw i32 [[VAL]], 1
-// CHECK: }
+// CK0: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]])
+// CK0: [[ADDR:%.+]] = alloca %class.C*,
+// CK0: store %class.C* [[ARG]], %class.C** [[ADDR]]
+// CK0: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
+// CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
+// CK0: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
+// CK0: {{.+}} = add nsw i32 [[VAL]], 1
+// CK0: }
+
+#endif
+
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+
+#ifdef CK1
+
+template <class T>
+class C {
+public:
+ T a;
+};
+
+#pragma omp declare mapper(id: C<int> s) map(s.a)
+
+// Synchronous version of mapper function.
+// CK1-LABEL: define {{.*}}i32 @".omp_mapper.C<int>.id{{.*}}(i64, i8*, i8*, i{{64|32}}, i64)
+// CK1-DAG: store i64 %0, i64* [[DIDADDR:%[^,]+]]
+// CK1-DAG: store i[[sz:64|32]] %3, i{{64|32}}* [[SIZEADDR:%[^,]+]]
+// CK1-DAG: store i64 %4, i64* [[TYPEADDR:%[^,]+]]
+// CK1-DAG: store i8* %1, i8** [[BPTRADDR:%[^,]+]]
+// CK1-DAG: store i8* %2, i8** [[VPTRADDR:%[^,]+]]
+// CK1-DAG: store i32 0, i32* %retval
+// CK1-DAG: [[SIZE:%.+]] = load i[[sz]], i[[sz]]* [[SIZEADDR]]
+// CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
+// CK1-DAG: [[DID:%.+]] = load i64, i64* [[DIDADDR]]
+// CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
+// CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i[[sz]] [[SIZE]]
+// CK1: [[ISARRAY:%.+]] = icmp sge i[[sz]] [[SIZE]], 1
+// CK1: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
+
+// CK1: [[INITEVALDEL]]
+// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK1: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
+// CK1: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK1: [[INIT]]
+// CK1: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 4
+// CK1-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[ISIZEADDR:[^,]+]]
+// CK1-DAG: [[ISIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[ISIZE:%.+]], i32 0, i32 0
+// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: store i64 [[ITYPE]], i64* [[ITYPEADDR:[^,]+]]
+// CK1-DAG: [[ITYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[ITYPE:%.+]], i32 0, i32 0
+// CK1: [[IRES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[ISIZEADDR]], i64* [[ITYPEADDR]], i8** null)
+// CK1: [[ISINITERR:%.+]] = icmp ne i32 [[IRES]], 0
+// CK1: br i1 [[ISINITERR]], label %[[INITERR:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK1: [[INITERR]]
+// CK1: store i32 [[IRES]], i32* %retval
+// CK1: br label %[[DONE:[^,]+]]
+
+// CK1: [[LHEAD]]
+// CK1: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
+// CK1: br i1 [[ISEMPTY]], label %[[DONE]], label %[[LBODY:[^,]+]]
+// CK1: [[LBODY]]
+// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
+// CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
+// CK1-DAG: [[AEND:%.+]] = getelementptr i32, i32* [[ABEGIN]], i32 1
+// CK1-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
+// CK1-DAG: [[AENDV:%.+]] = bitcast i32* [[AEND]] to i8*
+// CK1-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
+// CK1-DAG: [[AENDI:%.+]] = ptrtoint i8* [[AENDV]] to i64
+// CK1-DAG: [[CSIZE:%.+]] = sub i64 [[AENDI]], [[ABEGINI]]
+// CK1-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK1-32-DAG: [[CUSIZE32:%.+]] = trunc i64 [[CUSIZE]] to i32
+// CK1-64-DAG: store i[[sz]] [[CUSIZE]], i[[sz]]* [[SIZEADDR0:%[^,]+]]
+// CK1-32-DAG: store i[[sz]] [[CUSIZE32]], i[[sz]]* [[SIZEADDR0:%[^,]+]]
+// CK1-DAG: [[SIZEADDR0]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR:%[^,]+]], i32 0, i32 0
+// CK1-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR0BC:%[^,]+]]
+// CK1-DAG: [[BPTRADDR0BC]] = bitcast i8** [[BPTRADDR0:%.+]] to %class.C**
+// CK1-DAG: [[BPTRADDR0]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR:%[^,]+]], i32 0, i32 0
+// CK1-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR0BC:%[^,]+]]
+// CK1-DAG: [[PTRADDR0BC]] = bitcast i8** [[PTRADDR0:%.+]] to i32**
+// CK1-DAG: [[PTRADDR0]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR:%[^,]+]], i32 0, i32 0
+// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK1-DAG: [[ALLOC]]
+// CK1-DAG: br label %[[TYEND:[^,]+]]
+// CK1-DAG: [[ALLOCELSE]]
+// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK1-DAG: [[TO]]
+// CK1-DAG: br label %[[TYEND]]
+// CK1-DAG: [[TOELSE]]
+// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK1-DAG: [[FROM]]
+// CK1-DAG: br label %[[TYEND]]
+// CK1-DAG: [[TYEND]]
+// CK1-DAG: [[PHITYPE0:%.+]] = phi i64 [ 32, %[[ALLOC]] ], [ 32, %[[TO]] ], [ 32, %[[FROM]] ], [ 32, %[[TOELSE]] ]
+// CK1-DAG: store i64 [[PHITYPE0]], i64* [[TYPEADDR0:%[^,]+]]
+// CK1-DAG: [[TYPEADDR0]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR:%[^,]+]], i32 0, i32 0
+// CK1-DAG: [[BPTRADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR]], i32 0, i32 1
+// CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast i8** [[BPTRADDR1]] to %class.C**
+// CK1-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR1BC]]
+// CK1-DAG: [[PTRADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR]], i32 0, i32 1
+// CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i8** [[PTRADDR1]] to i32**
+// CK1-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR1BC]]
+// CK1-DAG: [[SIZEADDR1:%.+]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 1
+// CK1-DAG: store i[[sz]] 4, i[[sz]]* [[SIZEADDR1]]
+// CK1-DAG: [[TYPEADDR1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR]], i32 0, i32 1
+// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK1-DAG: [[ALLOC]]
+// CK1-DAG: br label %[[TYEND:[^,]+]]
+// CK1-DAG: [[ALLOCELSE]]
+// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK1-DAG: [[TO]]
+// CK1-DAG: br label %[[TYEND]]
+// CK1-DAG: [[TOELSE]]
+// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK1-DAG: [[FROM]]
+// CK1-DAG: br label %[[TYEND]]
+// CK1-DAG: [[TYEND]]
+// 0x1,000,000,000; 0x1,000,000,001; 0x1,000,000,002; 0x1,000,000,003
+// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ 281474976710656, %[[ALLOC]] ], [ 281474976710657, %[[TO]] ], [ 281474976710658, %[[FROM]] ], [ 281474976710659, %[[TOELSE]] ]
+// CK1-DAG: store i64 [[TYPE1]], i64* [[TYPEADDR1]]
+// CK1-DAG: [[ARGBPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR]], i32 0, i32 0
+// CK1-DAG: [[ARGPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR]], i32 0, i32 0
+// CK1-DAG: [[ARGSIZE:%.+]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 0
+// CK1-DAG: [[ARGTYPE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR]], i32 0, i32 0
+// CK1: [[RES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 2, i8** [[ARGBPTR]], i8** [[ARGPTR]], i[[sz]]* [[ARGSIZE]], i64* [[ARGTYPE]], i8** null)
+// CK1: [[ISERR:%.+]] = icmp ne i32 [[RES]], 0
+// CK1: br i1 [[ISERR]], label %[[LERR:[^,]+]], label %[[LCORRECT]]
+// CK1: [[LERR]]
+// CK1: store i32 [[RES]], i32* %retval
+// CK1: br label %[[DONE]]
+// CK1: [[LCORRECT]]
+// CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
+// CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
+// CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
+
+// CK1: [[LEXIT]]
+// CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
+// CK1: [[EVALDEL]]
+// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
+// CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
+// CK1: [[DEL]]
+// CK1: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 4
+// CK1-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[DSIZEADDR:%[^,]+]]
+// CK1-DAG: [[DSIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[DSIZE:%.+]], i32 0, i32 0
+// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: store i64 [[DTYPE]], i64* [[DTYPEADDR:%[^,]+]]
+// CK1-DAG: [[DTYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[DTYPE:%.+]], i32 0, i32 0
+// CK1: [[DRES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[DSIZEADDR]], i64* [[DTYPEADDR]], i8** null)
+// CK1: [[ISDELERR:%.+]] = icmp ne i32 [[DRES]], 0
+// CK1: br i1 [[ISDELERR]], label %[[DELERR:[^,]+]], label %[[DONE]]
+// CK1: [[DELERR]]
+// CK1: store i32 [[DRES]], i32* %retval
+// CK1: br label %[[DONE]]
+// CK1: [[DONE]]
+// CK1: [[RET:%.+]] = load i32, i32* %retval
+// CK1: ret i32 [[RET]]
+
+
+// Asynchronous version of mapper function.
+// CK1-LABEL: define {{.*}}i32 @".omp_mapper.C<int>.id{{.*}}nowait{{.*}}(i64, i8*, i8*, i{{64|32}}, i64)
+// CK1-DAG: store i64 %0, i64* [[DIDADDR:%[^,]+]]
+// CK1-DAG: store i[[sz:64|32]] %3, i{{64|32}}* [[SIZEADDR:%[^,]+]]
+// CK1-DAG: store i64 %4, i64* [[TYPEADDR:%[^,]+]]
+// CK1-DAG: store i8* %1, i8** [[BPTRADDR:%[^,]+]]
+// CK1-DAG: store i8* %2, i8** [[VPTRADDR:%[^,]+]]
+// CK1-DAG: store i32 0, i32* %retval
+// CK1-DAG: [[SIZE:%.+]] = load i[[sz]], i[[sz]]* [[SIZEADDR]]
+// CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
+// CK1-DAG: [[DID:%.+]] = load i64, i64* [[DIDADDR]]
+// CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
+// CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i[[sz]] [[SIZE]]
+// CK1: [[ISARRAY:%.+]] = icmp sge i[[sz]] [[SIZE]], 1
+// CK1: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
+
+// CK1: [[INITEVALDEL]]
+// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK1: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
+// CK1: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK1: [[INIT]]
+// CK1: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 4
+// CK1-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[ISIZEADDR:[^,]+]]
+// CK1-DAG: [[ISIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[ISIZE:%.+]], i32 0, i32 0
+// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: store i64 [[ITYPE]], i64* [[ITYPEADDR:[^,]+]]
+// CK1-DAG: [[ITYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[ITYPE:%.+]], i32 0, i32 0
+// CK1: [[IRES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[ISIZEADDR]], i64* [[ITYPEADDR]], i8** null)
+// CK1: [[ISINITERR:%.+]] = icmp ne i32 [[IRES]], 0
+// CK1: br i1 [[ISINITERR]], label %[[INITERR:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK1: [[INITERR]]
+// CK1: store i32 [[IRES]], i32* %retval
+// CK1: br label %[[DONE:[^,]+]]
+
+// CK1: [[LHEAD]]
+// CK1: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
+// CK1: br i1 [[ISEMPTY]], label %[[DONE]], label %[[LBODY:[^,]+]]
+// CK1: [[LBODY]]
+// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
+// CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
+// CK1-DAG: [[AEND:%.+]] = getelementptr i32, i32* [[ABEGIN]], i32 1
+// CK1-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
+// CK1-DAG: [[AENDV:%.+]] = bitcast i32* [[AEND]] to i8*
+// CK1-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
+// CK1-DAG: [[AENDI:%.+]] = ptrtoint i8* [[AENDV]] to i64
+// CK1-DAG: [[CSIZE:%.+]] = sub i64 [[AENDI]], [[ABEGINI]]
+// CK1-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK1-32-DAG: [[CUSIZE32:%.+]] = trunc i64 [[CUSIZE]] to i32
+// CK1-64-DAG: store i[[sz]] [[CUSIZE]], i[[sz]]* [[SIZEADDR0:%[^,]+]]
+// CK1-32-DAG: store i[[sz]] [[CUSIZE32]], i[[sz]]* [[SIZEADDR0:%[^,]+]]
+// CK1-DAG: [[SIZEADDR0]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR:%[^,]+]], i32 0, i32 0
+// CK1-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR0BC:%[^,]+]]
+// CK1-DAG: [[BPTRADDR0BC]] = bitcast i8** [[BPTRADDR0:%.+]] to %class.C**
+// CK1-DAG: [[BPTRADDR0]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR:%[^,]+]], i32 0, i32 0
+// CK1-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR0BC:%[^,]+]]
+// CK1-DAG: [[PTRADDR0BC]] = bitcast i8** [[PTRADDR0:%.+]] to i32**
+// CK1-DAG: [[PTRADDR0]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR:%[^,]+]], i32 0, i32 0
+// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK1-DAG: [[ALLOC]]
+// CK1-DAG: br label %[[TYEND:[^,]+]]
+// CK1-DAG: [[ALLOCELSE]]
+// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK1-DAG: [[TO]]
+// CK1-DAG: br label %[[TYEND]]
+// CK1-DAG: [[TOELSE]]
+// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK1-DAG: [[FROM]]
+// CK1-DAG: br label %[[TYEND]]
+// CK1-DAG: [[TYEND]]
+// CK1-DAG: [[PHITYPE0:%.+]] = phi i64 [ 32, %[[ALLOC]] ], [ 32, %[[TO]] ], [ 32, %[[FROM]] ], [ 32, %[[TOELSE]] ]
+// CK1-DAG: store i64 [[PHITYPE0]], i64* [[TYPEADDR0:%[^,]+]]
+// CK1-DAG: [[TYPEADDR0]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR:%[^,]+]], i32 0, i32 0
+// CK1-DAG: [[BPTRADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR]], i32 0, i32 1
+// CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast i8** [[BPTRADDR1]] to %class.C**
+// CK1-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR1BC]]
+// CK1-DAG: [[PTRADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR]], i32 0, i32 1
+// CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i8** [[PTRADDR1]] to i32**
+// CK1-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR1BC]]
+// CK1-DAG: [[SIZEADDR1:%.+]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 1
+// CK1-DAG: store i[[sz]] 4, i[[sz]]* [[SIZEADDR1]]
+// CK1-DAG: [[TYPEADDR1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR]], i32 0, i32 1
+// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK1-DAG: [[ALLOC]]
+// CK1-DAG: br label %[[TYEND:[^,]+]]
+// CK1-DAG: [[ALLOCELSE]]
+// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK1-DAG: [[TO]]
+// CK1-DAG: br label %[[TYEND]]
+// CK1-DAG: [[TOELSE]]
+// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK1-DAG: [[FROM]]
+// CK1-DAG: br label %[[TYEND]]
+// CK1-DAG: [[TYEND]]
+// 0x1,000,000,000; 0x1,000,000,001; 0x1,000,000,002; 0x1,000,000,003
+// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ 281474976710656, %[[ALLOC]] ], [ 281474976710657, %[[TO]] ], [ 281474976710658, %[[FROM]] ], [ 281474976710659, %[[TOELSE]] ]
+// CK1-DAG: store i64 [[TYPE1]], i64* [[TYPEADDR1]]
+// CK1-DAG: [[ARGBPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR]], i32 0, i32 0
+// CK1-DAG: [[ARGPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR]], i32 0, i32 0
+// CK1-DAG: [[ARGSIZE:%.+]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 0
+// CK1-DAG: [[ARGTYPE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR]], i32 0, i32 0
+// CK1: [[RES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 2, i8** [[ARGBPTR]], i8** [[ARGPTR]], i[[sz]]* [[ARGSIZE]], i64* [[ARGTYPE]], i8** null)
+// CK1: [[ISERR:%.+]] = icmp ne i32 [[RES]], 0
+// CK1: br i1 [[ISERR]], label %[[LERR:[^,]+]], label %[[LCORRECT]]
+// CK1: [[LERR]]
+// CK1: store i32 [[RES]], i32* %retval
+// CK1: br label %[[DONE]]
+// CK1: [[LCORRECT]]
+// CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
+// CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
+// CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
+
+// CK1: [[LEXIT]]
+// CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
+// CK1: [[EVALDEL]]
+// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
+// CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
+// CK1: [[DEL]]
+// CK1: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 4
+// CK1-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[DSIZEADDR:%[^,]+]]
+// CK1-DAG: [[DSIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[DSIZE:%.+]], i32 0, i32 0
+// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: store i64 [[DTYPE]], i64* [[DTYPEADDR:%[^,]+]]
+// CK1-DAG: [[DTYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[DTYPE:%.+]], i32 0, i32 0
+// CK1: [[DRES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[DSIZEADDR]], i64* [[DTYPEADDR]], i8** null)
+// CK1: [[ISDELERR:%.+]] = icmp ne i32 [[DRES]], 0
+// CK1: br i1 [[ISDELERR]], label %[[DELERR:[^,]+]], label %[[DONE]]
+// CK1: [[DELERR]]
+// CK1: store i32 [[DRES]], i32* %retval
+// CK1: br label %[[DONE]]
+// CK1: [[DONE]]
+// CK1: [[RET:%.+]] = load i32, i32* %retval
+// CK1: ret i32 [[RET]]
+
+#endif
#endif
Index: lib/CodeGen/ModuleBuilder.cpp
===================================================================
--- lib/CodeGen/ModuleBuilder.cpp
+++ lib/CodeGen/ModuleBuilder.cpp
@@ -232,6 +232,9 @@
if (auto *DRD = dyn_cast<OMPDeclareReductionDecl>(Member)) {
if (Ctx->DeclMustBeEmitted(DRD))
Builder->EmitGlobal(DRD);
+ } else if (auto *DMD = dyn_cast<OMPDeclareMapperDecl>(Member)) {
+ if (Ctx->DeclMustBeEmitted(DMD))
+ Builder->EmitGlobal(DMD);
}
}
}
Index: lib/CodeGen/CodeGenModule.h
===================================================================
--- lib/CodeGen/CodeGenModule.h
+++ lib/CodeGen/CodeGenModule.h
@@ -1255,8 +1255,7 @@
CodeGenFunction *CGF = nullptr);
/// Emit a code for declare mapper construct.
- void EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
- CodeGenFunction *CGF = nullptr);
+ void EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D);
/// Emit a code for requires directive.
/// \param D Requires declaration
Index: lib/CodeGen/CGOpenMPRuntime.h
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.h
+++ lib/CodeGen/CGOpenMPRuntime.h
@@ -345,6 +345,12 @@
SmallVector<const OMPDeclareReductionDecl *, 4>>
FunctionUDRMapTy;
FunctionUDRMapTy FunctionUDRMap;
+ /// Map from the user-defined mapper declaration to its corresponding
+ /// functions. The first one is the synchronous version, while the second one
+ /// is the asynchronous version.
+ llvm::DenseMap<const OMPDeclareMapperDecl *,
+ std::pair<llvm::Function *, llvm::Function *>>
+ UDMMap;
/// Type kmp_critical_name, originally defined as typedef kmp_int32
/// kmp_critical_name[8];
llvm::ArrayType *KmpCriticalNameTy;
@@ -787,6 +793,22 @@
virtual std::pair<llvm::Function *, llvm::Function *>
getUserDefinedReduction(const OMPDeclareReductionDecl *D);
+ /// Emit code for the user defined mapper construct.
+ virtual void emitUserDefinedMapper(const OMPDeclareMapperDecl *D);
+
+ /// Emit a function for a user defined mapper. Whether it is synchronous or
+ /// asynchronous depends on \p NoWait.
+ virtual llvm::Function *emitUDMapperFunc(const OMPDeclareMapperDecl *D,
+ bool NoWait);
+
+ /// Emit the array initialization or deletion portion for user-defined mapper
+ /// code generation.
+ virtual llvm::Value *
+ emitUDMapperArrayInitOrDel(CodeGenFunction &MapperCGF, llvm::Value *DeviceID,
+ llvm::Value *BasePtr, llvm::Value *Ptr,
+ llvm::Value *Size, llvm::Value *MapType,
+ CharUnits ElementSize, bool IsInit, bool NoWait);
+
/// Emits outlined function for the specified OpenMP parallel directive
/// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
/// kmp_int32 BoundID, struct context_vars*).
@@ -2069,6 +2091,9 @@
bool IsOffloadEntry,
const RegionCodeGenTy &CodeGen) override;
+ /// Emit code for the user defined mapper construct.
+ void emitUserDefinedMapper(const OMPDeclareMapperDecl *D);
+
/// Emit the target offloading code associated with \a D. The emitted
/// code attempts offloading the execution to the device, an the event of
/// a failure it executes the host version outlined in \a OutlinedFn.
Index: lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.cpp
+++ lib/CodeGen/CGOpenMPRuntime.cpp
@@ -719,6 +719,14 @@
// arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
// *arg_types);
OMPRTL__tgt_target_data_update_nowait,
+ // Call to void __tgt_target_data_mapper(int64_t device_id, int32_t arg_num,
+ // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types, void
+ // **arg_mapper_ptrs);
+ OMPRTL__tgt_target_data_mapper,
+ // Call to void __tgt_target_data_mapper_nowait(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mapper_ptrs);
+ OMPRTL__tgt_target_data_mapper_nowait,
};
/// A basic class for pre|post-action for advanced codegen sequence for OpenMP
@@ -2401,6 +2409,38 @@
RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_nowait");
break;
}
+ case OMPRTL__tgt_target_data_mapper: {
+ // Build void __tgt_target_data_mapper(int64_t device_id, int32_t arg_num,
+ // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types,
+ // void **arg_mapper_ptrs);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.SizeTy->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.IntTy, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_data_mapper_nowait: {
+ // Build void __tgt_target_data_mapper_nowait(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mapper_ptrs);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.SizeTy->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.IntTy, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_mapper_nowait");
+ break;
+ }
}
assert(RTLFn && "Unable to find OpenMP runtime function");
return RTLFn;
@@ -7063,8 +7103,11 @@
: IE(IE), VD(VD) {}
};
- /// Directive from where the map clauses were extracted.
- const OMPExecutableDirective &CurDir;
+ /// The target directive from where the mappable clauses were extracted. It
+ /// is either a executable directive or a user-defined mapper directive.
+ llvm::PointerUnion<const OMPExecutableDirective *,
+ const OMPDeclareMapperDecl *>
+ CurDir;
/// Function the directive is being generated for.
CodeGenFunction &CGF;
@@ -7737,7 +7780,7 @@
public:
MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
- : CurDir(Dir), CGF(CGF) {
+ : CurDir(&Dir), CGF(CGF) {
// Extract firstprivate clause information.
for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
for (const auto *D : C->varlists())
@@ -7749,6 +7792,10 @@
DevPointersMap[L.first].push_back(L.second);
}
+ /// Constructor for the declare mapper directive.
+ MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF)
+ : CurDir(&Dir), CGF(CGF) {}
+
/// Generate code for the combined entry if we have a partially mapped struct
/// and take care of the mapping flags of the arguments corresponding to
/// individual struct members.
@@ -7811,17 +7858,20 @@
};
// FIXME: MSVC 2013 seems to require this-> to find member CurDir.
- for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>())
+ assert(this->CurDir.is<const OMPExecutableDirective *>() &&
+ "Expect a executable directive");
+ const auto *CurExecDir = this->CurDir.get<const OMPExecutableDirective *>();
+ 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());
}
- for (const auto *C : this->CurDir.getClausesOfKind<OMPToClause>())
+ 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 : this->CurDir.getClausesOfKind<OMPFromClause>())
+ 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());
@@ -7836,9 +7886,8 @@
llvm::MapVector<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
DeferredInfo;
- // FIXME: MSVC 2013 seems to require this-> to find member CurDir.
for (const auto *C :
- this->CurDir.getClausesOfKind<OMPUseDevicePtrClause>()) {
+ CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
for (const auto &L : C->component_lists()) {
assert(!L.second.empty() && "Not expecting empty list of components!");
const ValueDecl *VD = L.second.back().getAssociatedDeclaration();
@@ -7964,6 +8013,80 @@
}
}
+ /// Generate all the base pointers, section pointers, sizes and map types for
+ /// the extracted map clauses of user-defined mapper.
+ void generateAllInfoForMapper(MapBaseValuesArrayTy &BasePointers,
+ MapValuesArrayTy &Pointers,
+ MapValuesArrayTy &Sizes,
+ MapFlagsArrayTy &Types) const {
+ // FIXME: MSVC 2013 seems to require this-> to find member CurDir.
+ assert(this->CurDir.is<const OMPDeclareMapperDecl *>() &&
+ "Expect a declare mapper directive");
+ const auto *CurMapperDir = this->CurDir.get<const OMPDeclareMapperDecl *>();
+ // 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.
+ llvm::MapVector<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
+
+ // Helper function to fill the information map for the different supported
+ // clauses.
+ auto &&InfoGen = [&Info](
+ const ValueDecl *D,
+ OMPClauseMappableExprCommon::MappableExprComponentListRef L,
+ OpenMPMapClauseKind MapType,
+ ArrayRef<OpenMPMapModifierKind> MapModifiers,
+ bool ReturnDevicePointer, bool IsImplicit) {
+ const ValueDecl *VD =
+ D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
+ Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
+ IsImplicit);
+ };
+
+ for (const auto *C : CurMapperDir->clauselists()) {
+ const auto *MC = cast<OMPMapClause>(C);
+ for (const auto &L : MC->component_lists()) {
+ InfoGen(L.first, L.second, MC->getMapType(), MC->getMapTypeModifiers(),
+ /*ReturnDevicePointer=*/false, MC->isImplicit());
+ }
+ }
+
+ for (const auto &M : Info) {
+ // We need to know when we generate information for the first component
+ // associated with a capture, because the mapping flags depend on it.
+ bool IsFirstComponentList = true;
+
+ // Temporary versions of arrays
+ MapBaseValuesArrayTy CurBasePointers;
+ MapValuesArrayTy CurPointers;
+ MapValuesArrayTy CurSizes;
+ MapFlagsArrayTy CurTypes;
+ StructRangeInfoTy PartialStruct;
+
+ for (const MapInfo &L : M.second) {
+ assert(!L.Components.empty() &&
+ "Not expecting declaration with no component lists.");
+ // FIXME: MSVC 2013 seems to require this-> to find the member method.
+ this->generateInfoForComponentList(
+ L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+ CurPointers, CurSizes, CurTypes, PartialStruct,
+ IsFirstComponentList, L.IsImplicit);
+ 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())
+ 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());
+ }
+ }
+
/// Emit capture info for lambdas for variables captured by reference.
void generateInfoForLambdaCaptures(
const ValueDecl *VD, llvm::Value *Arg, MapBaseValuesArrayTy &BasePointers,
@@ -8072,7 +8195,10 @@
OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool>;
SmallVector<MapData, 4> DeclComponentLists;
// FIXME: MSVC 2013 seems to require this-> to find member CurDir.
- for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
+ assert(this->CurDir.is<const OMPExecutableDirective *>() &&
+ "Expect a executable directive");
+ const auto *CurExecDir = this->CurDir.get<const OMPExecutableDirective *>();
+ for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
for (const auto &L : C->decl_component_lists(VD)) {
assert(L.first == VD &&
"We got information for the wrong declaration??");
@@ -8220,9 +8346,13 @@
MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes,
MapFlagsArrayTy &Types) const {
+ // FIXME: MSVC 2013 seems to require this-> to find member CurDir.
+ assert(this->CurDir.is<const OMPExecutableDirective *>() &&
+ "Expect a executable directive");
+ const auto *CurExecDir = this->CurDir.get<const OMPExecutableDirective *>();
// Map other list items in the map clause which are not captured variables
// but "declare target link" global variables.,
- for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
+ for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
for (const auto &L : C->component_lists()) {
if (!L.first)
continue;
@@ -8332,7 +8462,8 @@
MappableExprsHandler::MapValuesArrayTy &Pointers,
MappableExprsHandler::MapValuesArrayTy &Sizes,
MappableExprsHandler::MapFlagsArrayTy &MapTypes,
- CGOpenMPRuntime::TargetDataInfo &Info) {
+ CGOpenMPRuntime::TargetDataInfo &Info,
+ llvm::Value *MapperMapType = nullptr) {
CodeGenModule &CGM = CGF.CGM;
ASTContext &Ctx = CGF.getContext();
@@ -8350,6 +8481,9 @@
break;
}
+ // Indicate whether it is code generation within a user-defined mapper.
+ bool IsMapper = MapperMapType;
+
llvm::APInt PointerNumAP(32, Info.NumberOfPtrs, /*isSigned=*/true);
QualType PointerArrayType =
Ctx.getConstantArrayType(Ctx.VoidPtrTy, PointerNumAP, ArrayType::Normal,
@@ -8387,20 +8521,31 @@
Info.SizesArray = SizesArrayGbl;
}
- // The map types are always constant so we don't need to generate code to
- // fill arrays. Instead, we create an array constant.
- SmallVector<uint64_t, 4> Mapping(MapTypes.size(), 0);
- llvm::copy(MapTypes, Mapping.begin());
- llvm::Constant *MapTypesArrayInit =
- llvm::ConstantDataArray::get(CGF.Builder.getContext(), Mapping);
- std::string MaptypesName =
- CGM.getOpenMPRuntime().getName({"offload_maptypes"});
- auto *MapTypesArrayGbl = new llvm::GlobalVariable(
- CGM.getModule(), MapTypesArrayInit->getType(),
- /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
- MapTypesArrayInit, MaptypesName);
- MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
- Info.MapTypesArray = MapTypesArrayGbl;
+ if (IsMapper) {
+ // Allocate the memory space for map types because they are decided
+ // dynamically.
+ QualType MapArrayType = Ctx.getConstantArrayType(
+ Ctx.getIntTypeForBitwidth(/*DestWidth*/ 64, /*Signed*/ true),
+ PointerNumAP, ArrayType::Normal,
+ /*IndexTypeQuals=*/0);
+ Info.MapTypesArray =
+ CGF.CreateMemTemp(MapArrayType, ".offload_maptypes").getPointer();
+ } else {
+ // The map types are always constant so we don't need to generate code to
+ // fill arrays. Instead, we create an array constant.
+ SmallVector<uint64_t, 4> Mapping(MapTypes.size(), 0);
+ llvm::copy(MapTypes, Mapping.begin());
+ llvm::Constant *MapTypesArrayInit =
+ llvm::ConstantDataArray::get(CGF.Builder.getContext(), Mapping);
+ std::string MaptypesName =
+ CGM.getOpenMPRuntime().getName({"offload_maptypes"});
+ auto *MapTypesArrayGbl = new llvm::GlobalVariable(
+ CGM.getModule(), MapTypesArrayInit->getType(),
+ /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
+ MapTypesArrayInit, MaptypesName);
+ MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
+ Info.MapTypesArray = MapTypesArrayGbl;
+ }
for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
llvm::Value *BPVal = *BasePointers[I];
@@ -8436,9 +8581,77 @@
CGF.Builder.CreateIntCast(Sizes[I], CGM.SizeTy, /*isSigned=*/true),
SAddr);
}
+
+ if (IsMapper) {
+ // Combine the map type inherited from user-defined mapper with that
+ // specified in the program.
+ // [OpenMP 5.0], 1.2.6. map-type decay.
+ // | alloc | to | from | tofrom | release | delete
+ // ----------------------------------------------------------
+ // alloc | alloc | alloc | alloc | alloc | release | delete
+ // to | alloc | to | alloc | to | release | delete
+ // from | alloc | alloc | from | from | release | delete
+ // tofrom | alloc | to | from | tofrom | release | delete
+ llvm::Value *GEP = CGF.Builder.CreateConstInBoundsGEP2_32(
+ llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs),
+ Info.MapTypesArray, /*Idx0=*/0, /*Idx1=*/I);
+ llvm::Value *OriMapType = CGF.Builder.getInt64(MapTypes[I]);
+ llvm::Value *LeftToFrom = CGF.Builder.CreateAnd(
+ MapperMapType,
+ CGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_TO |
+ MappableExprsHandler::OMP_MAP_FROM));
+ llvm::BasicBlock *AllocBB = CGF.createBasicBlock("omp.type.alloc");
+ llvm::BasicBlock *AllocElseBB =
+ CGF.createBasicBlock("omp.type.alloc.else");
+ llvm::BasicBlock *ToBB = CGF.createBasicBlock("omp.type.to");
+ llvm::BasicBlock *ToElseBB = CGF.createBasicBlock("omp.type.to.else");
+ llvm::BasicBlock *FromBB = CGF.createBasicBlock("omp.type.from");
+ llvm::BasicBlock *EndBB = CGF.createBasicBlock("omp.type.end");
+ llvm::Value *IsAlloc = CGF.Builder.CreateIsNull(LeftToFrom);
+ CGF.Builder.CreateCondBr(IsAlloc, AllocBB, AllocElseBB);
+ // In case of alloc, clear OMP_MAP_TO and OMP_MAP_FROM.
+ CGF.EmitBlock(AllocBB);
+ llvm::Value *AllocMapType = CGF.Builder.CreateAnd(
+ OriMapType,
+ CGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO |
+ MappableExprsHandler::OMP_MAP_FROM)));
+ CGF.Builder.CreateBr(EndBB);
+ CGF.EmitBlock(AllocElseBB);
+ llvm::Value *IsTo = CGF.Builder.CreateICmpEQ(
+ LeftToFrom, CGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_TO));
+ CGF.Builder.CreateCondBr(IsTo, ToBB, ToElseBB);
+ // In case of to, clear OMP_MAP_FROM.
+ CGF.EmitBlock(ToBB);
+ llvm::Value *ToMapType = CGF.Builder.CreateAnd(
+ OriMapType,
+ CGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_FROM));
+ CGF.Builder.CreateBr(EndBB);
+ CGF.EmitBlock(ToElseBB);
+ llvm::Value *IsFrom = CGF.Builder.CreateICmpEQ(
+ LeftToFrom,
+ CGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_FROM));
+ CGF.Builder.CreateCondBr(IsFrom, FromBB, EndBB);
+ // In case of from, clear OMP_MAP_TO.
+ CGF.EmitBlock(FromBB);
+ llvm::Value *FromMapType = CGF.Builder.CreateAnd(
+ OriMapType,
+ CGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_TO));
+ // In case of tofrom, do nothing.
+ CGF.EmitBlock(EndBB);
+ llvm::PHINode *MapType =
+ CGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.maptype");
+ MapType->addIncoming(AllocMapType, AllocBB);
+ MapType->addIncoming(ToMapType, ToBB);
+ MapType->addIncoming(FromMapType, FromBB);
+ MapType->addIncoming(OriMapType, ToElseBB);
+ Address Addr(GEP, Ctx.getTypeAlignInChars(Ctx.getIntTypeForBitwidth(
+ /*DestWidth*/ 64, /*Signed*/ true)));
+ CGF.Builder.CreateStore(MapType, Addr);
+ }
}
}
}
+
/// Emit the arguments to be passed to the runtime library based on the
/// arrays of pointers, sizes and map types.
static void emitOffloadingArraysArgument(
@@ -8569,6 +8782,322 @@
return nullptr;
}
+/// Emit code for the user defined mapper construct.
+void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D) {
+ if (UDMMap.count(D) > 0)
+ return;
+ // Generate a synchronous mapper function.
+ llvm::Function *SyncFn = emitUDMapperFunc(D, /*NoWait=*/false);
+ // Generate an asynchronous mapper function.
+ llvm::Function *AsyncFn = emitUDMapperFunc(D, /*NoWait=*/true);
+ // Add the generated mapper functions to UDMMap.
+ UDMMap.try_emplace(D, SyncFn, AsyncFn);
+}
+
+/// Emit the user-defined mapper function. Whether it is synchronous or
+/// asynchronous depends on \p NoWait. The code generation follows the
+/// pattern in the example below.
+/// \code
+/// int .omp_mapper_<mapper_id>.(int64_t device_id, Ty *base_ptr, Ty *ptr,
+/// size_t size, int64_t maptype) {
+/// // Allocate space for an array section first.
+/// if (size > 1 && !maptype.IsDelete) {
+/// res = __tgt_target_data_mapper(device_id, /*arg_num*/1, &base_ptr, &ptr,
+/// size*sizeof(Ty), maptype, null);
+/// if (res != 0) // Data mapping failed.
+/// return res;
+/// }
+/// // Map members.
+/// for (unsigned i = 0; i < size; i++) {
+/// ...; // Prepare arguments of __tgt_target_data_mapper.
+/// res = __tgt_target_data_mapper(device_id, arg_num, arg_base, arg, size,
+/// maptype, mapper);
+/// if (res != 0) // Data mapping failed.
+/// return res;
+/// }
+/// // Delete the array section.
+/// if (size > 1 && maptype.IsDelete) {
+/// res = __tgt_target_data_mapper(device_id, /*arg_num*/1, &base_ptr, &ptr,
+/// size*sizeof(Ty), maptype, null);
+/// if (res != 0) // Data mapping failed.
+/// return res;
+/// }
+/// }
+/// \endcode
+llvm::Function *CGOpenMPRuntime::emitUDMapperFunc(const OMPDeclareMapperDecl *D,
+ bool NoWait) {
+ ASTContext &C = CGM.getContext();
+ QualType Ty = D->getType();
+ QualType PtrTy = C.getPointerType(Ty).withRestrict();
+ QualType SizeTy = C.getSizeType();
+ QualType Int64Ty = C.getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/true);
+ auto *MapperVarDecl =
+ cast<VarDecl>(cast<DeclRefExpr>(D->getMapperVarRef())->getDecl());
+ SourceLocation Loc = D->getLocation();
+ CharUnits ElementSize = C.getTypeSizeInChars(Ty);
+
+ // Prepare mapper function arguments and attributes.
+ ImplicitParamDecl DeviceIdArg(C, Int64Ty, ImplicitParamDecl::Other);
+ ImplicitParamDecl BasePtrArg(C, /*DC=*/nullptr, MapperVarDecl->getLocation(),
+ /*Id=*/nullptr, C.VoidPtrTy,
+ ImplicitParamDecl::Other);
+ ImplicitParamDecl PtrArg(C, /*DC=*/nullptr, MapperVarDecl->getLocation(),
+ /*Id=*/nullptr, C.VoidPtrTy,
+ ImplicitParamDecl::Other);
+ ImplicitParamDecl SizeArg(C, SizeTy, ImplicitParamDecl::Other);
+ ImplicitParamDecl MapTypeArg(C, Int64Ty, ImplicitParamDecl::Other);
+ FunctionArgList Args;
+ Args.push_back(&DeviceIdArg);
+ Args.push_back(&BasePtrArg);
+ Args.push_back(&PtrArg);
+ Args.push_back(&SizeArg);
+ Args.push_back(&MapTypeArg);
+ const CGFunctionInfo &FnInfo =
+ CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.IntTy, Args);
+ llvm::FunctionType *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
+ std::string Name = getName(
+ {"omp_mapper", Ty.getAsString(), D->getName(), NoWait ? "nowait." : ""});
+ std::replace(Name.begin(), Name.end(), ' ', '_');
+ auto *Fn = llvm::Function::Create(FnTy, llvm::GlobalValue::InternalLinkage,
+ Name, &CGM.getModule());
+ CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, FnInfo);
+ Fn->removeFnAttr(llvm::Attribute::OptimizeNone);
+ // Start the mapper function code generation.
+ CodeGenFunction MapperCGF(CGM);
+ MapperCGF.StartFunction(GlobalDecl(), C.IntTy, Fn, FnInfo, Args, Loc, Loc);
+ // Initiate the return value to 0, which represents success.
+ llvm::Value *SuccessRetVal = llvm::ConstantInt::getNullValue(CGM.IntTy);
+ MapperCGF.EmitStoreOfScalar(SuccessRetVal, MapperCGF.ReturnValue,
+ /*Volatile=*/false, C.IntTy);
+ // Compute the starting and end addreses of array elements.
+ llvm::Value *Size = MapperCGF.EmitLoadOfScalar(
+ MapperCGF.GetAddrOfLocalVar(&SizeArg), /*Volatile=*/false,
+ C.getPointerType(SizeTy), Loc);
+ llvm::Value *Ptr = MapperCGF.GetAddrOfLocalVar(&PtrArg).getPointer();
+ llvm::Value *PtrBegin = MapperCGF.Builder.CreateBitCast(
+ Ptr, CGM.getTypes().ConvertTypeForMem(C.getPointerType(PtrTy)));
+ llvm::Value *PtrEnd = MapperCGF.Builder.CreateGEP(PtrBegin, Size);
+ llvm::Value *NullMapperArrayArg =
+ llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
+ llvm::Value *MapType = MapperCGF.EmitLoadOfScalar(
+ MapperCGF.GetAddrOfLocalVar(&MapTypeArg), /*Volatile=*/false,
+ C.getPointerType(Int64Ty), Loc);
+ // Prepare some common arguments.
+ llvm::Value *DeviceID = MapperCGF.EmitLoadOfScalar(
+ MapperCGF.GetAddrOfLocalVar(&DeviceIdArg), /*Volatile=*/false,
+ C.getPointerType(Int64Ty), Loc);
+ llvm::Value *BasePtr = MapperCGF.GetAddrOfLocalVar(&BasePtrArg).getPointer();
+
+ // Evaluate if this is an array section.
+ llvm::BasicBlock *IsNotDeleteBB =
+ MapperCGF.createBasicBlock("omp.arrayinit.evaldelete");
+ llvm::BasicBlock *ArrayInitBB = MapperCGF.createBasicBlock("omp.arrayinit");
+ llvm::BasicBlock *HeadBB = MapperCGF.createBasicBlock("omp.arraymap.head");
+ llvm::Value *IsArray = MapperCGF.Builder.CreateICmpSGE(
+ Size, MapperCGF.Builder.getIntN(C.getTypeSize(SizeTy), 1),
+ "omp.arrayinit.isarray");
+ MapperCGF.Builder.CreateCondBr(IsArray, IsNotDeleteBB, HeadBB);
+ // Evaluate if we are going to delete this section.
+ MapperCGF.EmitBlock(IsNotDeleteBB);
+ llvm::Value *DeleteBit = MapperCGF.Builder.CreateAnd(
+ MapType,
+ MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE));
+ llvm::Value *IsNotDelete =
+ MapperCGF.Builder.CreateIsNull(DeleteBit, "omp.arrayinit.notdelete");
+ MapperCGF.Builder.CreateCondBr(IsNotDelete, ArrayInitBB, HeadBB);
+
+ // Allocate the space if this is an array.
+ MapperCGF.EmitBlock(ArrayInitBB);
+ llvm::Value *InitReturn =
+ emitUDMapperArrayInitOrDel(MapperCGF, DeviceID, BasePtr, Ptr, Size,
+ MapType, ElementSize, /*IsInit=*/true, NoWait);
+
+ // Jump to the function end if the return value indicates data mapping failed.
+ llvm::BasicBlock *InitErrorBB =
+ MapperCGF.createBasicBlock("omp.arrayinit.error");
+ llvm::BasicBlock *ExitBB = MapperCGF.createBasicBlock("omp.arraymap.exit");
+ llvm::BasicBlock *DoneBB = MapperCGF.createBasicBlock("omp.done");
+ llvm::Value *IsInitFailed = MapperCGF.Builder.CreateIsNotNull(InitReturn);
+ MapperCGF.Builder.CreateCondBr(IsInitFailed, InitErrorBB, HeadBB);
+ MapperCGF.EmitBlock(InitErrorBB);
+ MapperCGF.EmitStoreOfScalar(InitReturn, MapperCGF.ReturnValue,
+ /*Volatile=*/false, C.IntTy);
+ MapperCGF.Builder.CreateBr(DoneBB);
+
+ // Emit a for loop to iterate through SizeArg of elements and map all of them.
+
+ // Emit the loop header block.
+ MapperCGF.EmitBlock(HeadBB);
+ llvm::BasicBlock *BodyBB = MapperCGF.createBasicBlock("omp.arraymap.body");
+ llvm::BasicBlock *ErrorBB = MapperCGF.createBasicBlock("omp.arraymap.error");
+ llvm::BasicBlock *CorrectBB =
+ MapperCGF.createBasicBlock("omp.arraymap.correct");
+ // Evaluate whether the initial condition is satisfied.
+ llvm::Value *IsEmpty =
+ MapperCGF.Builder.CreateICmpEQ(PtrBegin, PtrEnd, "omp.arraymap.isempty");
+ MapperCGF.Builder.CreateCondBr(IsEmpty, DoneBB, BodyBB);
+ llvm::BasicBlock *EntryBB = MapperCGF.Builder.GetInsertBlock();
+
+ // Emit the loop body block.
+ MapperCGF.EmitBlock(BodyBB);
+ llvm::PHINode *PtrPHI = MapperCGF.Builder.CreatePHI(
+ PtrBegin->getType(), 2, "omp.arraymap.ptrcurrent");
+ PtrPHI->addIncoming(PtrBegin, EntryBB);
+ Address PtrCurrent =
+ Address(PtrPHI, MapperCGF.GetAddrOfLocalVar(&PtrArg)
+ .getAlignment()
+ .alignmentOfArrayElement(ElementSize));
+ // Privatize the declared variable of mapper to be the current array element.
+ CodeGenFunction::OMPPrivateScope Scope(MapperCGF);
+ Scope.addPrivate(MapperVarDecl, [&MapperCGF, PtrCurrent, PtrTy]() {
+ return MapperCGF
+ .EmitLoadOfPointerLValue(PtrCurrent, PtrTy->castAs<PointerType>())
+ .getAddress();
+ });
+ (void)Scope.Privatize();
+
+ // Get map clause information.
+ // Fill up the arrays with all the mapped variables.
+ MappableExprsHandler::MapBaseValuesArrayTy BasePointers;
+ MappableExprsHandler::MapValuesArrayTy Pointers;
+ MappableExprsHandler::MapValuesArrayTy Sizes;
+ MappableExprsHandler::MapFlagsArrayTy MapTypes;
+ MappableExprsHandler MEHandler(*D, MapperCGF);
+ MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes);
+ // Fill up the arrays and create the arguments.
+ TargetDataInfo Info;
+ emitOffloadingArrays(MapperCGF, BasePointers, Pointers, Sizes, MapTypes, Info,
+ MapType);
+ llvm::Value *BasePointersArrayArg = nullptr;
+ llvm::Value *PointersArrayArg = nullptr;
+ llvm::Value *SizesArrayArg = nullptr;
+ llvm::Value *MapTypesArrayArg = nullptr;
+ emitOffloadingArraysArgument(MapperCGF, BasePointersArrayArg,
+ PointersArrayArg, SizesArrayArg,
+ MapTypesArrayArg, Info);
+
+ // Call the runtime API __tgt_target_data_mapper(_nowait) to map data.
+ llvm::Value *PointerNum = MapperCGF.Builder.getInt32(Info.NumberOfPtrs);
+ llvm::Value *OffloadingArgs[] = {
+ DeviceID, PointerNum, BasePointersArrayArg, PointersArrayArg,
+ SizesArrayArg, MapTypesArrayArg, NullMapperArrayArg};
+ llvm::Value *Return = MapperCGF.EmitRuntimeCall(
+ createRuntimeFunction(NoWait ? OMPRTL__tgt_target_data_mapper_nowait
+ : OMPRTL__tgt_target_data_mapper),
+ OffloadingArgs);
+
+ // Break the loop if the return value indicates data mapping failed.
+ llvm::Value *IsFailed = MapperCGF.Builder.CreateIsNotNull(Return);
+ MapperCGF.Builder.CreateCondBr(IsFailed, ErrorBB, CorrectBB);
+ MapperCGF.EmitBlock(ErrorBB);
+ MapperCGF.EmitStoreOfScalar(Return, MapperCGF.ReturnValue,
+ /*Volatile=*/false, C.IntTy);
+ MapperCGF.Builder.CreateBr(DoneBB);
+
+ // Update the pointer to point to the next element that needs to be mapped,
+ // and check whether we have mapped all elements.
+ MapperCGF.EmitBlock(CorrectBB);
+ llvm::Value *PtrNext = MapperCGF.Builder.CreateConstGEP1_32(
+ PtrPHI, /*Idx0=*/1, "omp.arraymap.next");
+ PtrPHI->addIncoming(PtrNext, CorrectBB);
+ llvm::Value *IsDone =
+ MapperCGF.Builder.CreateICmpEQ(PtrNext, PtrEnd, "omp.arraymap.isdone");
+ MapperCGF.Builder.CreateCondBr(IsDone, ExitBB, BodyBB);
+
+ // Delete the array section if specified by the input map type.
+ MapperCGF.EmitBlock(ExitBB);
+ llvm::BasicBlock *IsDeleteBB =
+ MapperCGF.createBasicBlock("omp.arraydel.evaldelete");
+ llvm::BasicBlock *ArrayDelBB = MapperCGF.createBasicBlock("omp.arraydel");
+ MapperCGF.Builder.CreateCondBr(IsArray, IsDeleteBB, DoneBB);
+ MapperCGF.EmitBlock(IsDeleteBB);
+ llvm::Value *DB = MapperCGF.Builder.CreateAnd(
+ MapType,
+ MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE));
+ llvm::Value *IsDelete =
+ MapperCGF.Builder.CreateIsNotNull(DB, "omp.arraydel.isdelete");
+ MapperCGF.Builder.CreateCondBr(IsDelete, ArrayDelBB, DoneBB);
+ MapperCGF.EmitBlock(ArrayDelBB);
+ llvm::Value *DelReturn = emitUDMapperArrayInitOrDel(
+ MapperCGF, DeviceID, BasePtr, Ptr, Size, MapType, ElementSize,
+ /*IsInit=*/false, NoWait);
+
+ // Jump to the function end if the return value indicates data mapping failed.
+ llvm::BasicBlock *DelErrorBB =
+ MapperCGF.createBasicBlock("omp.arraydel.error");
+ llvm::Value *IsDelFailed = MapperCGF.Builder.CreateIsNotNull(DelReturn);
+ MapperCGF.Builder.CreateCondBr(IsDelFailed, DelErrorBB, DoneBB);
+ MapperCGF.EmitBlock(DelErrorBB);
+ MapperCGF.EmitStoreOfScalar(DelReturn, MapperCGF.ReturnValue,
+ /*Volatile=*/false, C.IntTy);
+
+ // Emit the function exit block.
+ MapperCGF.EmitBlock(DoneBB, /*IsFinished=*/true);
+ MapperCGF.FinishFunction();
+ return Fn;
+}
+
+// Emit the array initialization or deletion portion for user-defined mapper
+// code generation.
+llvm::Value *CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
+ CodeGenFunction &MapperCGF, llvm::Value *DeviceID, llvm::Value *BasePtr,
+ llvm::Value *Ptr, llvm::Value *Size, llvm::Value *MapType,
+ CharUnits ElementSize, bool IsInit, bool NoWait) {
+ ASTContext &C = CGM.getContext();
+ QualType Int64Ty = C.getIntTypeForBitwidth(/*DestWidth*/ 64, /*Signed*/ true);
+ std::string Prefix = IsInit ? ".init" : ".del";
+ // Prepare the size argument.
+ unsigned SizeTyWidth = C.getTypeSize(C.getSizeType());
+ llvm::Value *ArraySize = MapperCGF.Builder.CreateMul(
+ Size, MapperCGF.Builder.getIntN(SizeTyWidth, ElementSize.getQuantity()));
+ llvm::APInt PointerNumAP(32, 1, /*isSigned=*/true);
+ QualType SizeArrayType =
+ C.getConstantArrayType(C.getSizeType(), PointerNumAP, ArrayType::Normal,
+ /*IndexTypeQuals=*/0);
+ llvm::Value *SizesArrayStorage =
+ MapperCGF.CreateMemTemp(SizeArrayType, Prefix + ".offload_sizes")
+ .getPointer();
+ llvm::Value *SizesArrayArg = MapperCGF.Builder.CreateConstInBoundsGEP2_32(
+ llvm::ArrayType::get(CGM.SizeTy, 1), SizesArrayStorage, /*Idx0=*/0,
+ /*Idx1=*/0);
+ Address SizesArrayAddr(SizesArrayArg, C.getTypeAlignInChars(C.getSizeType()));
+ MapperCGF.EmitStoreOfScalar(ArraySize, SizesArrayAddr, /*Volatile=*/false,
+ C.getSizeType());
+ // Prepare the map type argument.
+ QualType MapArrayType =
+ C.getConstantArrayType(Int64Ty, PointerNumAP, ArrayType::Normal,
+ /*IndexTypeQuals=*/0);
+ llvm::Value *MapTypeArrayStorage =
+ MapperCGF.CreateMemTemp(MapArrayType, Prefix + ".offload_maptypes")
+ .getPointer();
+ llvm::Value *MapTypeArg = MapperCGF.Builder.CreateConstInBoundsGEP2_32(
+ llvm::ArrayType::get(CGM.Int64Ty, 1), MapTypeArrayStorage, /*Idx0=*/0,
+ /*Idx1=*/0);
+ Address MapTypeArrayAddr(MapTypeArg, C.getTypeAlignInChars(Int64Ty));
+ // Remove OMP_MAP_TO and OMP_MAP_FROM from the map type, so that it achieves
+ // memory allocation/deletion purpose only.
+ llvm::Value *ArrayMapType = MapperCGF.Builder.CreateAnd(
+ MapType,
+ MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO |
+ MappableExprsHandler::OMP_MAP_FROM)));
+ MapperCGF.EmitStoreOfScalar(ArrayMapType, MapTypeArrayAddr,
+ /*Volatile=*/false, Int64Ty);
+ llvm::Value *NullMapperArrayArg =
+ llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
+ llvm::Value *OffloadingArgs[] = {
+ DeviceID,
+ /*arg_num*/ MapperCGF.Builder.getInt32(1),
+ BasePtr,
+ Ptr,
+ SizesArrayArg,
+ MapTypeArg,
+ NullMapperArrayArg};
+ return MapperCGF.EmitRuntimeCall(
+ createRuntimeFunction(NoWait ? OMPRTL__tgt_target_data_mapper_nowait
+ : OMPRTL__tgt_target_data_mapper),
+ OffloadingArgs);
+}
+
void CGOpenMPRuntime::emitTargetNumIterationsCall(
CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *Device,
const llvm::function_ref<llvm::Value *(
@@ -10586,6 +11115,10 @@
llvm_unreachable("Not supported in SIMD-only mode");
}
+void CGOpenMPSIMDRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D) {
+ llvm_unreachable("Not supported in SIMD-only mode");
+}
+
void CGOpenMPSIMDRuntime::emitTargetCall(CodeGenFunction &CGF,
const OMPExecutableDirective &D,
llvm::Function *OutlinedFn,
Index: lib/CodeGen/CGDecl.cpp
===================================================================
--- lib/CodeGen/CGDecl.cpp
+++ lib/CodeGen/CGDecl.cpp
@@ -144,7 +144,7 @@
return CGM.EmitOMPDeclareReduction(cast<OMPDeclareReductionDecl>(&D), this);
case Decl::OMPDeclareMapper:
- return CGM.EmitOMPDeclareMapper(cast<OMPDeclareMapperDecl>(&D), this);
+ return CGM.EmitOMPDeclareMapper(cast<OMPDeclareMapperDecl>(&D));
case Decl::Typedef: // typedef int X;
case Decl::TypeAlias: { // using X = int; [C++0x]
@@ -2495,11 +2495,11 @@
getOpenMPRuntime().emitUserDefinedReduction(CGF, D);
}
-void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
- CodeGenFunction *CGF) {
- if (!LangOpts.OpenMP || (!LangOpts.EmitAllDecls && !D->isUsed()))
+void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D) {
+ if (!LangOpts.OpenMP || LangOpts.OpenMPSimd ||
+ (!LangOpts.EmitAllDecls && !D->isUsed()))
return;
- // FIXME: need to implement mapper code generation
+ getOpenMPRuntime().emitUserDefinedMapper(D);
}
void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) {
Index: lib/AST/ASTContext.cpp
===================================================================
--- lib/AST/ASTContext.cpp
+++ lib/AST/ASTContext.cpp
@@ -9793,7 +9793,7 @@
return !D->getDeclContext()->isDependentContext();
else if (isa<OMPAllocateDecl>(D))
return !D->getDeclContext()->isDependentContext();
- else if (isa<OMPDeclareReductionDecl>(D))
+ else if (isa<OMPDeclareReductionDecl>(D) || isa<OMPDeclareMapperDecl>(D))
return !D->getDeclContext()->isDependentContext();
else if (isa<ImportDecl>(D))
return true;
Index: include/clang/AST/GlobalDecl.h
===================================================================
--- include/clang/AST/GlobalDecl.h
+++ include/clang/AST/GlobalDecl.h
@@ -59,6 +59,7 @@
GlobalDecl(const CapturedDecl *D) { Init(D); }
GlobalDecl(const ObjCMethodDecl *D) { Init(D); }
GlobalDecl(const OMPDeclareReductionDecl *D) { Init(D); }
+ GlobalDecl(const OMPDeclareMapperDecl *D) { Init(D); }
GlobalDecl(const CXXConstructorDecl *D, CXXCtorType Type) : Value(D, Type) {}
GlobalDecl(const CXXDestructorDecl *D, CXXDtorType Type) : Value(D, Type) {}
GlobalDecl(const VarDecl *D, DynamicInitKind StubKind)
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
