r267812 - [OpenMP] Code generation for target enter data directive
Samuel Antao via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 27 16:07:29 PDT 2016
Author: sfantao
Date: Wed Apr 27 18:07:29 2016
New Revision: 267812
URL: http://llvm.org/viewvc/llvm-project?rev=267812&view=rev
Log:
[OpenMP] Code generation for target enter data directive
Summary: This patch adds support for the target enter data directive code generation.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: cfe-commits, fraggamuffin, caomhin
Differential Revision: http://reviews.llvm.org/D17368
Added:
cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp
Modified:
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=267812&r1=267811&r2=267812&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Wed Apr 27 18:07:29 2016
@@ -5839,3 +5839,67 @@ void CGOpenMPRuntime::emitTargetDataCall
EndThenRCG(CGF);
}
}
+
+void CGOpenMPRuntime::emitTargetEnterDataCall(CodeGenFunction &CGF,
+ const OMPExecutableDirective &D,
+ const Expr *IfCond,
+ const Expr *Device) {
+ if (!CGF.HaveInsertPoint())
+ return;
+
+ // Generate the code for the opening of the data environment.
+ auto &&ThenGen = [&D, &CGF, Device](CodeGenFunction &CGF, PrePostActionTy &) {
+ // Fill up the arrays with all the mapped variables.
+ MappableExprsHandler::MapValuesArrayTy BasePointers;
+ MappableExprsHandler::MapValuesArrayTy Pointers;
+ MappableExprsHandler::MapValuesArrayTy Sizes;
+ MappableExprsHandler::MapFlagsArrayTy MapTypes;
+
+ // Get map clause information.
+ MappableExprsHandler MCHandler(D, CGF);
+ MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+
+ llvm::Value *BasePointersArrayArg = nullptr;
+ llvm::Value *PointersArrayArg = nullptr;
+ llvm::Value *SizesArrayArg = nullptr;
+ llvm::Value *MapTypesArrayArg = nullptr;
+
+ // Fill up the arrays and create the arguments.
+ emitOffloadingArrays(CGF, BasePointersArrayArg, PointersArrayArg,
+ SizesArrayArg, MapTypesArrayArg, BasePointers,
+ Pointers, Sizes, MapTypes);
+ emitOffloadingArraysArgument(
+ CGF, BasePointersArrayArg, PointersArrayArg, SizesArrayArg,
+ MapTypesArrayArg, BasePointersArrayArg, PointersArrayArg, SizesArrayArg,
+ MapTypesArrayArg, BasePointers.size());
+
+ // Emit device ID if any.
+ llvm::Value *DeviceID = nullptr;
+ if (Device)
+ DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
+ CGF.Int32Ty, /*isSigned=*/true);
+ else
+ DeviceID = CGF.Builder.getInt32(OMP_DEVICEID_UNDEF);
+
+ // Emit the number of elements in the offloading arrays.
+ auto *PointerNum = CGF.Builder.getInt32(BasePointers.size());
+
+ llvm::Value *OffloadingArgs[] = {
+ DeviceID, PointerNum, BasePointersArrayArg,
+ PointersArrayArg, SizesArrayArg, MapTypesArrayArg};
+ auto &RT = CGF.CGM.getOpenMPRuntime();
+ CGF.EmitRuntimeCall(RT.createRuntimeFunction(OMPRTL__tgt_target_data_begin),
+ OffloadingArgs);
+ };
+
+ // In the event we get an if clause, we don't have to take any action on the
+ // else side.
+ auto &&ElseGen = [](CodeGenFunction &CGF, PrePostActionTy &) {};
+
+ if (IfCond) {
+ emitOMPIfClause(CGF, IfCond, ThenGen, ElseGen);
+ } else {
+ RegionCodeGenTy ThenGenRCG(ThenGen);
+ ThenGenRCG(CGF);
+ }
+}
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h?rev=267812&r1=267811&r2=267812&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h Wed Apr 27 18:07:29 2016
@@ -1051,6 +1051,16 @@ public:
const OMPExecutableDirective &D,
const Expr *IfCond, const Expr *Device,
const RegionCodeGenTy &CodeGen);
+
+ /// \brief Emit the target data mapping code associated with \a D.
+ /// \param D Directive to emit.
+ /// \param IfCond Expression evaluated in if clause associated with the target
+ /// directive, or null if no if clause is used.
+ /// \param Device Expression evaluated in device clause associated with the
+ /// target directive, or null if no device clause is used.
+ virtual void emitTargetEnterDataCall(CodeGenFunction &CGF,
+ const OMPExecutableDirective &D,
+ const Expr *IfCond, const Expr *Device);
};
} // namespace CodeGen
Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=267812&r1=267811&r2=267812&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Wed Apr 27 18:07:29 2016
@@ -3272,7 +3272,22 @@ void CodeGenFunction::EmitOMPTargetDataD
void CodeGenFunction::EmitOMPTargetEnterDataDirective(
const OMPTargetEnterDataDirective &S) {
- // TODO: codegen for target enter data.
+ // If we don't have target devices, don't bother emitting the data mapping
+ // code.
+ if (CGM.getLangOpts().OMPTargetTriples.empty())
+ return;
+
+ // Check if we have any if clause associated with the directive.
+ const Expr *IfCond = nullptr;
+ if (auto *C = S.getSingleClause<OMPIfClause>())
+ IfCond = C->getCondition();
+
+ // Check if we have any device clause associated with the directive.
+ const Expr *Device = nullptr;
+ if (auto *C = S.getSingleClause<OMPDeviceClause>())
+ Device = C->getDevice();
+
+ CGM.getOpenMPRuntime().emitTargetEnterDataCall(*this, S, IfCond, Device);
}
void CodeGenFunction::EmitOMPTargetExitDataDirective(
Added: cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp?rev=267812&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp (added)
+++ cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp Wed Apr 27 18:07:29 2016
@@ -0,0 +1,221 @@
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
+// RUN: %clang_cc1 -DCK1 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK1 --check-prefix CK1-64
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
+// RUN: %clang_cc1 -DCK1 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK1 --check-prefix CK1-32
+#ifdef CK1
+
+// CK1: [[ST:%.+]] = type { i32, double* }
+template <typename T>
+struct ST {
+ T a;
+ double *b;
+};
+
+ST<int> gb;
+double gc[100];
+
+// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] zeroinitializer
+
+// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 1]
+
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 5]
+
+// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 1, i32 97]
+
+// CK1-LABEL: _Z3fooi
+void foo(int arg) {
+ int la;
+ float lb[arg];
+
+ // Region 00
+ // CK1-DAG: call void @__tgt_target_data_begin(i32 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1-DAG: [[DEV]] = load i32, i32* %{{[^,]+}},
+ // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: store i8* bitcast ([100 x double]* @gc to i8*), i8** [[BP0]]
+ // CK1-DAG: store i8* bitcast ([100 x double]* @gc to i8*), i8** [[P0]]
+
+ // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ // CK1-NOT: __tgt_target_data_end
+ #pragma omp target enter data if(1+3-5) device(arg) map(alloc: gc)
+ {++arg;}
+
+ // Region 01
+ // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ #pragma omp target enter data map(to: la) if(1+3-4)
+ {++arg;}
+
+ // Region 02
+ // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
+ // CK1: [[IFTHEN]]
+ // CK1-DAG: call void @__tgt_target_data_begin(i32 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+ // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+ // CK1-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+ // CK1-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
+ // CK1-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8*
+ // CK1: br label %[[IFEND:[^,]+]]
+
+ // CK1: [[IFELSE]]
+ // CK1: br label %[[IFEND]]
+ // CK1: [[IFEND]]
+ // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ // CK1-NOT: __tgt_target_data_end
+ #pragma omp target enter data map(to: arg) if(arg) device(4)
+ {++arg;}
+
+ // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ {++arg;}
+
+ // Region 03
+ // CK1-DAG: call void @__tgt_target_data_begin(i32 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
+ // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+ // CK1-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+ // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+ // CK1-DAG: [[CBPVAL0]] = bitcast float* [[VAR0:%.+]] to i8*
+ // CK1-DAG: [[CPVAL0]] = bitcast float* [[VAR0]] to i8*
+ // CK1-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
+ // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ // CK1-NOT: __tgt_target_data_end
+ #pragma omp target enter data map(always, to: lb)
+ {++arg;}
+
+ // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ {++arg;}
+
+ // Region 04
+ // CK1-DAG: call void @__tgt_target_data_begin(i32 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}})
+ // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK1-DAG: store i8* bitcast ([[ST]]* @gb to i8*), i8** [[BP0]]
+ // CK1-DAG: store i8* bitcast (double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1) to i8*), i8** [[P0]]
+
+
+ // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK1-DAG: store i8* bitcast (double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1) to i8*), i8** [[BP1]]
+ // CK1-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+ // CK1-DAG: [[CPVAL1]] = bitcast double* [[SEC1:%.+]] to i8*
+ // CK1-DAG: [[SEC1]] = getelementptr inbounds {{.+}}double* [[SEC11:%[^,]+]], i{{.+}} 0
+ // CK1-DAG: [[SEC11]] = load double*, double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1),
+
+ // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ // CK1-NOT: __tgt_target_data_end
+ #pragma omp target enter data map(to: gb.b[:3])
+ {++arg;}
+}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
+// RUN: %clang_cc1 -DCK2 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK2 --check-prefix CK2-64
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
+// RUN: %clang_cc1 -DCK2 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK2 --check-prefix CK2-32
+#ifdef CK2
+
+// CK2: [[ST:%.+]] = type { i32, double* }
+template <typename T>
+struct ST {
+ T a;
+ double *b;
+
+ T foo(T arg) {
+ // Region 00
+ #pragma omp target enter data map(always, to: b[1:3]) if(a>123) device(arg)
+ {arg++;}
+ return arg;
+ }
+};
+
+// CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 5, i32 101]
+
+// CK2-LABEL: _Z3bari
+int bar(int arg){
+ ST<int> A;
+ return A.foo(arg);
+}
+
+// Region 00
+// CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
+// CK2: [[IFTHEN]]
+// CK2-DAG: call void @__tgt_target_data_begin(i32 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+// CK2-DAG: [[DEV]] = load i32, i32* %{{[^,]+}},
+// CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK2-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK2-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK2-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
+// CK2-DAG: [[CPVAL0]] = bitcast double** [[SEC0:%[^,]+]] to i8*
+// CK2-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[ST]]* [[VAR0]], i32 0, i32 1
+
+
+// CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK2-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK2-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK2-DAG: [[CBPVAL1]] = bitcast double** [[SEC0]] to i8*
+// CK2-DAG: [[CPVAL1]] = bitcast double* [[SEC1:%[^,]+]] to i8*
+// CK2-DAG: [[SEC1]] = getelementptr inbounds {{.*}}double* [[SEC11:%[^,]+]], i{{.+}} 1
+// CK2-DAG: [[SEC11]] = load double*, double** [[SEC111:%[^,]+]],
+// CK2-DAG: [[SEC111]] = getelementptr inbounds {{.*}}[[ST]]* [[VAR0]], i32 0, i32 1
+
+// CK2: br label %[[IFEND:[^,]+]]
+
+// CK2: [[IFELSE]]
+// CK2: br label %[[IFEND]]
+// CK2: [[IFEND]]
+// CK2: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+// CK2-NOT: __tgt_target_data_end
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
+// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
+// RUN: %clang_cc1 -DCK3 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32
+// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32
+#ifdef CK3
+
+// CK3-LABEL: no_target_devices
+void no_target_devices(int arg) {
+ // CK3-NOT: tgt_target_data_begin
+ // CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+ // CK3-NOT: tgt_target_data_end
+ // CK3: ret
+ #pragma omp target enter data map(to: arg) if(arg) device(4)
+ {++arg;}
+}
+#endif
+#endif
More information about the cfe-commits
mailing list