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