r310379 - Revert "[OPENMP][DEBUG] Set proper address space info if required by target."
David Blaikie via cfe-commits
cfe-commits at lists.llvm.org
Mon Aug 14 12:08:59 PDT 2017
It's helpful to mention why a patch is reverted in the commit message that
reverts. Thought for next time! :)
On Tue, Aug 8, 2017 at 9:46 AM Alexey Bataev via cfe-commits <
cfe-commits at lists.llvm.org> wrote:
> 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]]*
> -
>
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20170814/83c788ea/attachment-0001.html>
More information about the cfe-commits
mailing list