r310379 - Revert "[OPENMP][DEBUG] Set proper address space info if required by target."
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Tue Aug 8 09:45:36 PDT 2017
Author: abataev
Date: Tue Aug 8 09:45:36 2017
New Revision: 310379
URL: http://llvm.org/viewvc/llvm-project?rev=310379&view=rev
Log:
Revert "[OPENMP][DEBUG] Set proper address space info if required by target."
This reverts commit r310377.
Removed:
cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp
Modified:
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
cfe/trunk/lib/Sema/SemaExpr.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Tue Aug 8 09:45:36 2017
@@ -2685,14 +2685,6 @@ def OMPCaptureNoInit : InheritableAttr {
let Documentation = [Undocumented];
}
-def OMPCaptureKind : Attr {
- // This attribute has no spellings as it is only ever created implicitly.
- let Spellings = [];
- let SemaHandler = 0;
- let Args = [UnsignedArgument<"CaptureKind">];
- let Documentation = [Undocumented];
-}
-
def OMPDeclareSimdDecl : Attr {
let Spellings = [Pragma<"omp", "declare simd">];
let Subjects = SubjectList<[Function]>;
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Tue Aug 8 09:45:36 2017
@@ -8527,11 +8527,6 @@ public:
/// is performed.
bool isOpenMPPrivateDecl(ValueDecl *D, unsigned Level);
- /// Sets OpenMP capture kind (OMPC_private, OMPC_firstprivate, OMPC_map etc.)
- /// for \p FD based on DSA for the provided corresponding captured declaration
- /// \p D.
- void setOpenMPCaptureKind(FieldDecl *FD, ValueDecl *D, unsigned Level);
-
/// \brief Check if the specified variable is captured by 'target' directive.
/// \param Level Relative level of nested OpenMP construct for that the check
/// is performed.
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h Tue Aug 8 09:45:36 2017
@@ -1325,32 +1325,6 @@ public:
virtual void emitDoacrossOrdered(CodeGenFunction &CGF,
const OMPDependClause *C);
- /// Translates the native parameter of outlined function if this is required
- /// for target.
- /// \param FD Field decl from captured record for the paramater.
- /// \param NativeParam Parameter itself.
- virtual const VarDecl *translateParameter(const FieldDecl *FD,
- const VarDecl *NativeParam) const {
- return NativeParam;
- }
-
- typedef llvm::function_ref<void(CodeGenFunction &, const VarDecl *, Address)>
- MappingFnType;
- /// Maps the native argument to the address of the corresponding
- /// target-specific argument.
- /// \param FD Field decl from captured record for the paramater.
- /// \param NativeParam Parameter itself.
- /// \param TargetParam Corresponding target-specific parameter.
- /// \param MapFn Function that maps the native parameter to the address of the
- /// target-specific.
- virtual void mapParameterAddress(CodeGenFunction &CGF, const FieldDecl *FD,
- const VarDecl *NativeParam,
- const VarDecl *TargetParam,
- const MappingFnType) const {
- assert(NativeParam == TargetParam &&
- "native and target args must be the same");
- }
-
/// Emits call of the outlined function with the provided arguments,
/// translating these arguments to correct target-specific arguments.
virtual void
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp Tue Aug 8 09:45:36 2017
@@ -2238,81 +2238,3 @@ void CGOpenMPRuntimeNVPTX::emitReduction
CGF.EmitBranch(DefaultBB);
CGF.EmitBlock(DefaultBB, /*IsFinished=*/true);
}
-
-const VarDecl *
-CGOpenMPRuntimeNVPTX::translateParameter(const FieldDecl *FD,
- const VarDecl *NativeParam) const {
- if (!NativeParam->getType()->isReferenceType())
- return NativeParam;
- QualType ArgType = NativeParam->getType();
- QualifierCollector QC;
- const Type *NonQualTy = QC.strip(ArgType);
- QualType PointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
- if (const auto *Attr = FD->getAttr<OMPCaptureKindAttr>()) {
- if (Attr->getCaptureKind() == OMPC_map) {
- PointeeTy = CGM.getContext().getAddrSpaceQualType(PointeeTy,
- LangAS::opencl_global);
- }
- }
- ArgType = CGM.getContext().getPointerType(PointeeTy);
- QC.addRestrict();
- enum { NVPTX_local_addr = 5 };
- QC.addAddressSpace(NVPTX_local_addr);
- ArgType = QC.apply(CGM.getContext(), ArgType);
- return ImplicitParamDecl::Create(
- CGM.getContext(), /*DC=*/nullptr, NativeParam->getLocation(),
- NativeParam->getIdentifier(), ArgType, ImplicitParamDecl::Other);
-}
-
-void CGOpenMPRuntimeNVPTX::mapParameterAddress(
- CodeGenFunction &CGF, const FieldDecl *FD, const VarDecl *NativeParam,
- const VarDecl *TargetParam,
- const CGOpenMPRuntime::MappingFnType MapFn) const {
- assert(NativeParam != TargetParam &&
- NativeParam->getType()->isReferenceType() &&
- "Native arg must not be the same as target arg.");
- Address LocalAddr = CGF.GetAddrOfLocalVar(TargetParam);
- QualType NativeParamType = NativeParam->getType();
- QualifierCollector QC;
- const Type *NonQualTy = QC.strip(NativeParamType);
- QualType NativePointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
- unsigned NativePointeeAddrSpace =
- NativePointeeTy.getQualifiers().getAddressSpace();
- QualType TargetPointeeTy = TargetParam->getType()->getPointeeType();
- llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(
- LocalAddr, /*Volatile=*/false, TargetPointeeTy, SourceLocation());
- // First cast to generic.
- TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
- TargetAddr, TargetAddr->getType()->getPointerElementType()->getPointerTo(
- /*AddrSpace=*/0));
- // Cast from generic to native address space.
- TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
- TargetAddr, TargetAddr->getType()->getPointerElementType()->getPointerTo(
- NativePointeeAddrSpace));
- Address NativeParamAddr = CGF.CreateMemTemp(NativeParamType);
- CGF.EmitStoreOfScalar(TargetAddr, NativeParamAddr, /*Volatile=*/false,
- NativeParam->getType());
- MapFn(CGF, NativeParam, NativeParamAddr);
-}
-
-void CGOpenMPRuntimeNVPTX::emitOutlinedFunctionCall(
- CodeGenFunction &CGF, llvm::Value *OutlinedFn,
- ArrayRef<llvm::Value *> Args) const {
- SmallVector<llvm::Value *, 4> TargetArgs;
- auto *FnType =
- cast<llvm::FunctionType>(OutlinedFn->getType()->getPointerElementType());
- for (unsigned I = 0, E = Args.size(); I < E; ++I) {
- llvm::Type *TargetType = FnType->getParamType(I);
- llvm::Value *NativeArg = Args[I];
- if (!TargetType->isPointerTy()) {
- TargetArgs.emplace_back(NativeArg);
- continue;
- }
- llvm::Value *TargetArg = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
- NativeArg, NativeArg->getType()->getPointerElementType()->getPointerTo(
- /*AddrSpace=*/0));
- TargetArgs.emplace_back(
- CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(TargetArg, TargetType));
- }
- CGOpenMPRuntime::emitOutlinedFunctionCall(CGF, OutlinedFn, TargetArgs);
-}
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h Tue Aug 8 09:45:36 2017
@@ -268,31 +268,6 @@ public:
/// \return Specified function.
llvm::Constant *createNVPTXRuntimeFunction(unsigned Function);
- /// Translates the native parameter of outlined function if this is required
- /// for target.
- /// \param FD Field decl from captured record for the paramater.
- /// \param NativeParam Parameter itself.
- const VarDecl *translateParameter(const FieldDecl *FD,
- const VarDecl *NativeParam) const override;
-
- /// Maps the native argument to the address of the corresponding
- /// target-specific argument.
- /// \param FD Field decl from captured record for the paramater.
- /// \param NativeParam Parameter itself.
- /// \param TargetParam Corresponding target-specific parameter.
- /// \param MapFn Function that maps the native parameter to the address of the
- /// target-specific.
- void mapParameterAddress(CodeGenFunction &CGF, const FieldDecl *FD,
- const VarDecl *NativeParam,
- const VarDecl *TargetParam,
- const MappingFnType MapFn) const override;
-
- /// Emits call of the outlined function with the provided arguments,
- /// translating these arguments to correct target-specific arguments.
- void emitOutlinedFunctionCall(
- CodeGenFunction &CGF, llvm::Value *OutlinedFn,
- ArrayRef<llvm::Value *> Args = llvm::None) const override;
-
/// Target codegen is specialized based on two programming models: the
/// 'generic' fork-join model of OpenMP, and a more GPU efficient 'spmd'
/// model for constructs like 'target parallel' that support it.
Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Tue Aug 8 09:45:36 2017
@@ -246,27 +246,24 @@ namespace {
const CapturedStmt *S = nullptr;
/// true if cast to/from UIntPtr is required for variables captured by
/// value.
- const bool UIntPtrCastRequired = true;
+ bool UIntPtrCastRequired = true;
/// true if only casted argumefnts must be registered as local args or VLA
/// sizes.
- const bool RegisterCastedArgsOnly = false;
+ bool RegisterCastedArgsOnly = false;
/// Name of the generated function.
- const StringRef FunctionName;
- /// Function that maps given variable declaration to the specified address.
- const CGOpenMPRuntime::MappingFnType MapFn;
+ StringRef FunctionName;
explicit FunctionOptions(const CapturedStmt *S, bool UIntPtrCastRequired,
bool RegisterCastedArgsOnly,
- StringRef FunctionName,
- const CGOpenMPRuntime::MappingFnType MapFn)
+ StringRef FunctionName)
: S(S), UIntPtrCastRequired(UIntPtrCastRequired),
RegisterCastedArgsOnly(UIntPtrCastRequired && RegisterCastedArgsOnly),
- FunctionName(FunctionName), MapFn(MapFn) {}
+ FunctionName(FunctionName) {}
};
}
static std::pair<llvm::Function *, bool> emitOutlinedFunctionPrologue(
CodeGenFunction &CGF, FunctionArgList &Args,
- llvm::MapVector<const Decl *, std::pair<const VarDecl *, Address>>
+ llvm::DenseMap<const Decl *, std::pair<const VarDecl *, Address>>
&LocalAddrs,
llvm::DenseMap<const Decl *, std::pair<const Expr *, llvm::Value *>>
&VLASizes,
@@ -279,13 +276,9 @@ static std::pair<llvm::Function *, bool>
// Build the argument list.
CodeGenModule &CGM = CGF.CGM;
ASTContext &Ctx = CGM.getContext();
- FunctionArgList TargetArgs;
bool HasUIntPtrArgs = false;
Args.append(CD->param_begin(),
std::next(CD->param_begin(), CD->getContextParamPosition()));
- TargetArgs.append(
- CD->param_begin(),
- std::next(CD->param_begin(), CD->getContextParamPosition()));
auto I = FO.S->captures().begin();
for (auto *FD : RD->fields()) {
QualType ArgType = FD->getType();
@@ -315,28 +308,19 @@ static std::pair<llvm::Function *, bool>
}
if (ArgType->isVariablyModifiedType())
ArgType = getCanonicalParamType(Ctx, ArgType.getNonReferenceType());
- auto *Arg =
- ImplicitParamDecl::Create(Ctx, /*DC=*/nullptr, FD->getLocation(), II,
- ArgType, ImplicitParamDecl::Other);
- Args.emplace_back(Arg);
- // Do not cast arguments if we emit function with non-original types.
- TargetArgs.emplace_back(
- FO.UIntPtrCastRequired
- ? Arg
- : CGM.getOpenMPRuntime().translateParameter(FD, Arg));
+ Args.push_back(ImplicitParamDecl::Create(Ctx, /*DC=*/nullptr,
+ FD->getLocation(), II, ArgType,
+ ImplicitParamDecl::Other));
++I;
}
Args.append(
std::next(CD->param_begin(), CD->getContextParamPosition() + 1),
CD->param_end());
- TargetArgs.append(
- std::next(CD->param_begin(), CD->getContextParamPosition() + 1),
- CD->param_end());
// Create the function declaration.
FunctionType::ExtInfo ExtInfo;
const CGFunctionInfo &FuncInfo =
- CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, TargetArgs);
+ CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, Args);
llvm::FunctionType *FuncLLVMTy = CGM.getTypes().GetFunctionType(FuncInfo);
llvm::Function *F =
@@ -347,21 +331,16 @@ static std::pair<llvm::Function *, bool>
F->setDoesNotThrow();
// Generate the function.
- CGF.StartFunction(CD, Ctx.VoidTy, F, FuncInfo, TargetArgs, CD->getLocation(),
+ CGF.StartFunction(CD, Ctx.VoidTy, F, FuncInfo, Args, CD->getLocation(),
CD->getBody()->getLocStart());
unsigned Cnt = CD->getContextParamPosition();
I = FO.S->captures().begin();
for (auto *FD : RD->fields()) {
- // Do not map arguments if we emit function with non-original types.
- if (!FO.UIntPtrCastRequired && Args[Cnt] != TargetArgs[Cnt]) {
- CGM.getOpenMPRuntime().mapParameterAddress(CGF, FD, Args[Cnt],
- TargetArgs[Cnt], FO.MapFn);
- }
- Address LocalAddr = CGF.GetAddrOfLocalVar(Args[Cnt]);
// If we are capturing a pointer by copy we don't need to do anything, just
// use the value that we get from the arguments.
if (I->capturesVariableByCopy() && FD->getType()->isAnyPointerType()) {
const VarDecl *CurVD = I->getCapturedVar();
+ Address LocalAddr = CGF.GetAddrOfLocalVar(Args[Cnt]);
// If the variable is a reference we need to materialize it here.
if (CurVD->getType()->isReferenceType()) {
Address RefAddr = CGF.CreateMemTemp(
@@ -378,8 +357,8 @@ static std::pair<llvm::Function *, bool>
}
LValueBaseInfo BaseInfo(AlignmentSource::Decl, false);
- LValue ArgLVal =
- CGF.MakeAddrLValue(LocalAddr, Args[Cnt]->getType(), BaseInfo);
+ LValue ArgLVal = CGF.MakeAddrLValue(CGF.GetAddrOfLocalVar(Args[Cnt]),
+ Args[Cnt]->getType(), BaseInfo);
if (FD->hasCapturedVLAType()) {
if (FO.UIntPtrCastRequired) {
ArgLVal = CGF.MakeAddrLValue(castValueFromUintptr(CGF, FD->getType(),
@@ -447,19 +426,10 @@ CodeGenFunction::GenerateOpenMPCapturedS
getDebugInfo() &&
CGM.getCodeGenOpts().getDebugInfo() >= codegenoptions::LimitedDebugInfo;
FunctionArgList Args;
- llvm::MapVector<const Decl *, std::pair<const VarDecl *, Address>> LocalAddrs;
+ llvm::DenseMap<const Decl *, std::pair<const VarDecl *, Address>> LocalAddrs;
llvm::DenseMap<const Decl *, std::pair<const Expr *, llvm::Value *>> VLASizes;
- FunctionOptions FO(
- &S, !NeedWrapperFunction, /*RegisterCastedArgsOnly=*/false,
- CapturedStmtInfo->getHelperName(),
- [NeedWrapperFunction](CodeGenFunction &CGF, const VarDecl *VD,
- Address Addr) {
- if (!NeedWrapperFunction) {
- llvm_unreachable("Function should not be called if wrapper function "
- "is not required.");
- }
- CGF.setAddrOfLocalVar(VD, Addr);
- });
+ FunctionOptions FO(&S, !NeedWrapperFunction, /*RegisterCastedArgsOnly=*/false,
+ CapturedStmtInfo->getHelperName());
llvm::Function *F;
bool HasUIntPtrArgs;
std::tie(F, HasUIntPtrArgs) = emitOutlinedFunctionPrologue(
@@ -482,10 +452,7 @@ CodeGenFunction::GenerateOpenMPCapturedS
llvm::raw_svector_ostream Out(Buffer);
Out << "__nondebug_wrapper_" << CapturedStmtInfo->getHelperName();
FunctionOptions WrapperFO(&S, /*UIntPtrCastRequired=*/true,
- /*RegisterCastedArgsOnly=*/true, Out.str(),
- [](CodeGenFunction &, const VarDecl *, Address) {
- llvm_unreachable("Function should not be called");
- });
+ /*RegisterCastedArgsOnly=*/true, Out.str());
CodeGenFunction WrapperCGF(CGM, /*suppressNewContext=*/true);
WrapperCGF.disableDebugInfo();
Args.clear();
Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Tue Aug 8 09:45:36 2017
@@ -14013,8 +14013,6 @@ static bool captureInCapturedRegion(Capt
Field->setImplicit(true);
Field->setAccess(AS_private);
RD->addDecl(Field);
- if (S.getLangOpts().OpenMP && RSI->CapRegionKind == CR_OpenMP)
- S.setOpenMPCaptureKind(Field, Var, RSI->OpenMPLevel);
CopyExpr = new (S.Context) DeclRefExpr(Var, RefersToCapturedVariable,
DeclRefType, VK_LValue, Loc);
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Tue Aug 8 09:45:36 2017
@@ -1327,39 +1327,6 @@ bool Sema::isOpenMPPrivateDecl(ValueDecl
DSAStack->isTaskgroupReductionRef(D, Level));
}
-void Sema::setOpenMPCaptureKind(FieldDecl *FD, ValueDecl *D, unsigned Level) {
- assert(LangOpts.OpenMP && "OpenMP is not allowed");
- D = getCanonicalDecl(D);
- OpenMPClauseKind OMPC = OMPC_unknown;
- for (unsigned I = DSAStack->getNestingLevel() + 1; I > Level; --I) {
- const unsigned NewLevel = I - 1;
- if (DSAStack->hasExplicitDSA(D,
- [&OMPC](const OpenMPClauseKind K) {
- if (isOpenMPPrivate(K)) {
- OMPC = K;
- return true;
- }
- return false;
- },
- NewLevel))
- break;
- if (DSAStack->checkMappableExprComponentListsForDeclAtLevel(
- D, NewLevel,
- [](OMPClauseMappableExprCommon::MappableExprComponentListRef,
- OpenMPClauseKind) { return true; })) {
- OMPC = OMPC_map;
- break;
- }
- if (DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective,
- NewLevel)) {
- OMPC = OMPC_firstprivate;
- break;
- }
- }
- if (OMPC != OMPC_unknown)
- FD->addAttr(OMPCaptureKindAttr::CreateImplicit(Context, OMPC));
-}
-
bool Sema::isOpenMPTargetCapturedDecl(ValueDecl *D, unsigned Level) {
assert(LangOpts.OpenMP && "OpenMP is not allowed");
// Return true if the current level is no longer enclosed in a target region.
Modified: cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp?rev=310379&r1=310378&r2=310379&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp Tue Aug 8 09:45:36 2017
@@ -1,14 +1,15 @@
+
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
-// RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
-template <typename tx, typename ty>
-struct TT {
+template<typename tx, typename ty>
+struct TT{
tx X;
ty Y;
};
@@ -22,32 +23,29 @@ int foo(int n, double *ptr) {
float b[10];
double c[5][10];
TT<long long, char> d;
-
-#pragma omp target firstprivate(a) map(tofrom \
- : b)
+
+ #pragma omp target firstprivate(a)
{
- b[a] = a;
}
-
- // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}([10 x float] addrspace(1)* noalias [[B_IN:%.+]], i{{[0-9]+}} [[A_IN:%.+]])
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK-NOT: alloca i{{[0-9]+}},
- // TCHECK-64: call void @llvm.dbg.declare(metadata [10 x float] addrspace(1)** %{{.+}}, metadata !{{[0-9]+}}, metadata ![[LOCAL:[0-9]+]])
// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
- // TCHECK: ret void
+ // TCHECK: ret void
-#pragma omp target firstprivate(aa, b, c, d)
+#pragma omp target firstprivate(aa,b,c,d)
{
aa += 1;
b[2] = 1.0;
c[1][2] = 1.0;
d.X = 1;
- d.Y = 1;
+ d.Y = 1;
}
-
+
// make sure that firstprivate variables are generated in all cases and that we use those instances for operations inside the
// target region
- // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}(i{{[0-9]+}}{{.*}} [[A2_IN:%.+]], [10 x float]*{{.*}} [[B_IN:%.+]], [5 x [10 x double]]*{{.*}} [[C_IN:%.+]], [[TT]]*{{.*}} [[D_IN:%.+]])
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A2_IN:%.+]], [10 x float]* {{.+}} [[B_IN:%.+]], [5 x [10 x double]]* {{.+}} [[C_IN:%.+]], [[TT]]* {{.+}} [[D_IN:%.+]])
// TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[B_ADDR:%.+]] = alloca [10 x float]*,
// TCHECK: [[C_ADDR:%.+]] = alloca [5 x [10 x double]]*,
@@ -60,12 +58,10 @@ int foo(int n, double *ptr) {
// TCHECK: store [10 x float]* [[B_IN]], [10 x float]** [[B_ADDR]],
// TCHECK: store [5 x [10 x double]]* [[C_IN]], [5 x [10 x double]]** [[C_ADDR]],
// TCHECK: store [[TT]]* [[D_IN]], [[TT]]** [[D_ADDR]],
+ // TCHECK: [[CONV_A2ADDR:%.+]] = bitcast i{{[0-9]+}}* [[A2_ADDR]] to i{{[0-9]+}}*
// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** [[B_ADDR]],
- // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** %
// TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[C_ADDR]],
- // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** %
// TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** [[D_ADDR]],
- // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** %
// firstprivate(aa): a_priv = a_in
@@ -78,15 +74,16 @@ int foo(int n, double *ptr) {
// TCHECK: [[C_PRIV_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_PRIV]] to i8*
// TCHECK: [[C_IN_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8*
// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_PRIV_BCAST]], i8* [[C_IN_BCAST]],{{.+}})
-
+
// firstprivate(d)
// TCHECK: [[D_PRIV_BCAST:%.+]] = bitcast [[TT]]* [[D_PRIV]] to i8*
// TCHECK: [[D_IN_BCAST:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8*
// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[D_PRIV_BCAST]], i8* [[D_IN_BCAST]],{{.+}})
- // TCHECK: load i16, i16* [[A2_ADDR]],
+ // TCHECK: load i16, i16* [[CONV_A2ADDR]],
-#pragma omp target firstprivate(ptr)
+
+ #pragma omp target firstprivate(ptr)
{
ptr[0]++;
}
@@ -101,12 +98,13 @@ int foo(int n, double *ptr) {
return a;
}
-template <typename tx>
+
+template<typename tx>
tx ftemplate(int n) {
tx a = 0;
tx b[10];
-#pragma omp target firstprivate(a, b)
+#pragma omp target firstprivate(a,b)
{
a += 1;
b[2] += 1;
@@ -115,12 +113,13 @@ tx ftemplate(int n) {
return a;
}
-static int fstatic(int n) {
+static
+int fstatic(int n) {
int a = 0;
char aaa = 0;
int b[10];
-#pragma omp target firstprivate(a, aaa, b)
+#pragma omp target firstprivate(a,aaa,b)
{
a += 1;
aaa += 1;
@@ -130,7 +129,7 @@ static int fstatic(int n) {
return a;
}
-// TCHECK: define {{.*}}void @__omp_offloading_{{.+}}(i{{[0-9]+}}{{.*}} [[A_IN:%.+]], i{{[0-9]+}}{{.*}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
@@ -139,8 +138,9 @@ static int fstatic(int n) {
// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
// TCHECK: store i{{[0-9]+}} [[A3_IN]], i{{[0-9]+}}* [[A3_ADDR]],
// TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK-64: [[A_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[A3_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A3_ADDR]] to i8*
// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
-// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** %
// firstprivate(a): a_priv = a_in
@@ -158,8 +158,8 @@ static int fstatic(int n) {
struct S1 {
double a;
- int r1(int n) {
- int b = n + 1;
+ int r1(int n){
+ int b = n+1;
#pragma omp target firstprivate(b)
{
@@ -169,7 +169,7 @@ struct S1 {
return (int)b;
}
- // TCHECK: define internal void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]])
+ // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]])
// TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
// TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK-NOT: alloca i{{[0-9]+}},
@@ -185,7 +185,9 @@ struct S1 {
// TCHECK: ret void
};
-int bar(int n, double *ptr) {
+
+
+int bar(int n, double *ptr){
int a = 0;
a += foo(n, ptr);
S1 S;
@@ -198,15 +200,15 @@ int bar(int n, double *ptr) {
// template
-// TCHECK: define internal void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
// TCHECK-NOT: alloca i{{[0-9]+}},
// TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
// TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK-64: [[A_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
-// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** %
// firstprivate(a)
// TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}*
Removed: cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp?rev=310378&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_parallel_debug_codegen.cpp (removed)
@@ -1,97 +0,0 @@
-// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited | FileCheck %s
-// expected-no-diagnostics
-
-int main() {
- /* int(*b)[a]; */
- /* int *(**c)[a]; */
- int a;
- int b[10][10];
- int c[10][10][10];
-#pragma omp target parallel firstprivate(a, b) map(tofrom \
- : c)
- {
- int &f = c[1][1][1];
- int &g = a;
- int &h = b[1][1];
- int d = 15;
- a = 5;
- b[0][a] = 10;
- c[0][0][a] = 11;
- b[0][a] = c[0][0][a];
- }
-#pragma omp target parallel firstprivate(a) map(tofrom \
- : c, b)
- {
- int &f = c[1][1][1];
- int &g = a;
- int &h = b[1][1];
- int d = 15;
- a = 5;
- b[0][a] = 10;
- c[0][0][a] = 11;
- b[0][a] = c[0][0][a];
- }
-#pragma omp target parallel map(tofrom \
- : a, c, b)
- {
- int &f = c[1][1][1];
- int &g = a;
- int &h = b[1][1];
- int d = 15;
- a = 5;
- b[0][a] = 10;
- c[0][0][a] = 11;
- b[0][a] = c[0][0][a];
- }
- return 0;
-}
-
-// CHECK: define internal void @__omp_offloading{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
-// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* {{[^)]+}})
-
-// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* noalias{{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
-
-// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
-// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^)]+}})
-
-// CHECK: define void @__nondebug_wrapper___omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
-// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^)]+}})
-
-// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* noalias {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias {{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
-// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
-// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* {{[^)]+}})
-
-// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
-// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
-
-// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
-// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
-// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^)]+}})
-
-// CHECK: define void @__nondebug_wrapper___omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
-// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
-// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^)]+}})
-
-// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* noalias {{[^,]+}}, i32 addrspace(1)* noalias {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias {{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
-// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
-// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
-// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
-// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
-// CHECK: call void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^)]+}})
-
-// CHECK: define internal void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 addrspace(1)* noalias{{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^)]+}})
-// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
-// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
-// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
-
More information about the cfe-commits
mailing list