[OPENMP]Remove const firstprivate allocation as a variable in a constant space.

[ABataev]

Summary

Current implementation is not compatible with asynchronous target
regions, need to remove it.

[ABataev]

Rebase

[vzakhari]

LGTM

[jhuber6]

This revision causes a bug with generic regions. Firstprivate constants now aren't mapped properly into an internal parallel region and will just be zero. For example if I run this code I will see different values for the constants inside and outside the parallel region.

#include <complex>
#include <cstdio>

void foo(const std::complex<double> X, std::complex<double> Y) {
#pragma omp target firstprivate(X) map(from:Y)
  {
    printf ("Outside parallel: %f + %fi\n", X.real(), X.imag());
#pragma omp parallel firstprivate(X)
    {
      printf ("Inside parallel: %f + %fi\n", X.real(), X.imag());
      Y = X;
    }
  }
}

int main() {
  std::complex<double> X = {1.0, 1.0};
  std::complex<double> Y;
  foo(X, Y);
}

Outside parallel: 1.000000 + 1.000000i
Inside parallel: 0.000000 + 0.000000i

I'm assuming they aren't being properly copied because they were originally global constants that had visibility everywhere.

[ABataev]

In D105375#2864733, @jhuber6 wrote:

This revision causes a bug with generic regions. Firstprivate constants now aren't mapped properly into an internal parallel region and will just be zero. For example if I run this code I will see different values for the constants inside and outside the parallel region.

#include <complex>
#include <cstdio>

void foo(const std::complex<double> X, std::complex<double> Y) {
#pragma omp target firstprivate(X) map(from:Y)
  {
    printf ("Outside parallel: %f + %fi\n", X.real(), X.imag());
#pragma omp parallel firstprivate(X)
    {
      printf ("Inside parallel: %f + %fi\n", X.real(), X.imag());
      Y = X;
    }
  }
}

int main() {
  std::complex<double> X = {1.0, 1.0};
  std::complex<double> Y;
  foo(X, Y);
}

Outside parallel: 1.000000 + 1.000000i
Inside parallel: 0.000000 + 0.000000i

I'm assuming they aren't being properly copied because they were originally global constants that had visibility everywhere.

Hi, thanks for the report. I think I know the reason, I prepared a patch D105647. Could you check if it fixes the issue?

clang/lib/CodeGen/CGOpenMPRuntime.h

private:
   /// TaskFunction.
   /// \param Data Additional data for task generation like tiednsee, final
   /// state, list of privates etc.
   TaskResultTy emitTaskInit(CodeGenFunction &CGF, SourceLocation Loc,
                             const OMPExecutableDirective &D,
                             llvm::Function *TaskFunction, QualType SharedsTy,
                             Address Shareds, const OMPTaskDataTy &Data);
 
-  /// Returns default address space for the constant firstprivates, 0 by
-  /// default.
-  virtual unsigned getDefaultFirstprivateAddressSpace() const { return 0; }
-
   /// Emit code that pushes the trip count of loops associated with constructs
   /// 'target teams distribute' and 'teams distribute parallel for'.
   /// \param SizeEmitter Emits the int64 value for the number of iterations of
   /// the associated loop.
   void emitTargetNumIterationsCall(
       CodeGenFunction &CGF, const OMPExecutableDirective &D,
       llvm::Value *DeviceID,
       llvm::function_ref<llvm::Value *(CodeGenFunction &CGF,
public:
   /// \param GD Variable declaration to emit.
   virtual bool emitTargetGlobalVariable(GlobalDecl GD);
 
   /// Checks if the provided global decl \a GD is a declare target variable and
   /// registers it when emitting code for the host.
   virtual void registerTargetGlobalVariable(const VarDecl *VD,
                                             llvm::Constant *Addr);
 
-  /// Registers provided target firstprivate variable as global on the
-  /// target.
-  llvm::Constant *registerTargetFirstprivateCopy(CodeGenFunction &CGF,
-                                                 const VarDecl *VD);
-
   /// Emit the global \a GD if it is meaningful for the target. Returns
   /// if it was emitted successfully.
   /// \param GD Global to scan.
   virtual bool emitTargetGlobal(GlobalDecl GD);
 
   /// Creates and returns a registration function for when at least one
   /// requires directives was used in the current module.
   llvm::Function *emitRequiresDirectiveRegFun();

clang/lib/CodeGen/CGOpenMPRuntime.cpp

   MappableExprsHandler::OpenMPOffloadMappingFlags
   getMapModifiersForPrivateClauses(const CapturedStmt::Capture &Cap) const {
     assert(Cap.capturesVariable() && "Expected capture by reference only!");
 
     // A first private variable captured by reference will use only the
     // 'private ptr' and 'map to' flag. Return the right flags if the captured
     // declaration is known as first-private in this handler.
     if (FirstPrivateDecls.count(Cap.getCapturedVar())) {
-      if (Cap.getCapturedVar()->getType().isConstant(CGF.getContext()) &&
-          Cap.getCaptureKind() == CapturedStmt::VCK_ByRef)
-        return MappableExprsHandler::OMP_MAP_ALWAYS |
-               MappableExprsHandler::OMP_MAP_TO;
       if (Cap.getCapturedVar()->getType()->isAnyPointerType())
         return MappableExprsHandler::OMP_MAP_TO |
                MappableExprsHandler::OMP_MAP_PTR_AND_OBJ;
       return MappableExprsHandler::OMP_MAP_PRIVATE |
              MappableExprsHandler::OMP_MAP_TO;
     }
     return MappableExprsHandler::OMP_MAP_TO |
            MappableExprsHandler::OMP_MAP_FROM;
if (CI.capturesThis()) {
       CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
           CGF.getTypeSize(ElementType), CGF.Int64Ty, /*isSigned=*/true));
       // The default map type for a scalar/complex type is 'to' because by
       // default the value doesn't have to be retrieved. For an aggregate
       // type, the default is 'tofrom'.
       CombinedInfo.Types.push_back(getMapModifiersForPrivateClauses(CI));
       const VarDecl *VD = CI.getCapturedVar();
       auto I = FirstPrivateDecls.find(VD);
-      if (I != FirstPrivateDecls.end() &&
-          VD->getType().isConstant(CGF.getContext())) {
-        llvm::Constant *Addr =
-            CGF.CGM.getOpenMPRuntime().registerTargetFirstprivateCopy(CGF, VD);
-        // Copy the value of the original variable to the new global copy.
-        CGF.Builder.CreateMemCpy(
-            CGF.MakeNaturalAlignAddrLValue(Addr, ElementType).getAddress(CGF),
-            Address(CV, CGF.getContext().getTypeAlignInChars(ElementType)),
-            CombinedInfo.Sizes.back(), /*IsVolatile=*/false);
-        // Use new global variable as the base pointers.
-        CombinedInfo.Exprs.push_back(VD->getCanonicalDecl());
-        CombinedInfo.BasePointers.push_back(Addr);
-        CombinedInfo.Pointers.push_back(Addr);
-      } else {
-        CombinedInfo.Exprs.push_back(VD->getCanonicalDecl());
-        CombinedInfo.BasePointers.push_back(CV);
-        if (I != FirstPrivateDecls.end() && ElementType->isAnyPointerType()) {
-          Address PtrAddr = CGF.EmitLoadOfReference(CGF.MakeAddrLValue(
-              CV, ElementType, CGF.getContext().getDeclAlign(VD),
-              AlignmentSource::Decl));
-          CombinedInfo.Pointers.push_back(PtrAddr.getPointer());
-        } else {
-          CombinedInfo.Pointers.push_back(CV);
-        }
-      }
       if (I != FirstPrivateDecls.end())
         IsImplicit = I->getSecond();
     }
     // Every default map produces a single argument which is a target parameter.
     CombinedInfo.Types.back() |= OMP_MAP_TARGET_PARAM;
 
     // Add flag stating this is an implicit map.
     if (IsImplicit)
if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
       (*Res == OMPDeclareTargetDeclAttr::MT_To &&
        HasRequiresUnifiedSharedMemory)) {
     DeferredGlobalVariables.insert(cast<VarDecl>(GD.getDecl()));
     return true;
   }
   return false;
 }
 
-llvm::Constant *
-CGOpenMPRuntime::registerTargetFirstprivateCopy(CodeGenFunction &CGF,
-                                                const VarDecl *VD) {
-  assert(VD->getType().isConstant(CGM.getContext()) &&
-         "Expected constant variable.");
-  StringRef VarName;
-  llvm::Constant *Addr;
-  llvm::GlobalValue::LinkageTypes Linkage;
-  QualType Ty = VD->getType();
-  SmallString<128> Buffer;
-  {
-    unsigned DeviceID;
-    unsigned FileID;
-    unsigned Line;
-    getTargetEntryUniqueInfo(CGM.getContext(), VD->getLocation(), DeviceID,
-                             FileID, Line);
-    llvm::raw_svector_ostream OS(Buffer);
-    OS << "__omp_offloading_firstprivate_" << llvm::format("_%x", DeviceID)
-       << llvm::format("_%x_", FileID);
-    if (CGM.getLangOpts().CPlusPlus) {
-      CGM.getCXXABI().getMangleContext().mangleTypeName(VD->getType(), OS);
-      OS << "_";
-    }
-    OS << VD->getName() << "_l" << Line;
-    VarName = OS.str();
-  }
-  Linkage = llvm::GlobalValue::InternalLinkage;
-  Addr =
-      getOrCreateInternalVariable(CGM.getTypes().ConvertTypeForMem(Ty), VarName,
-                                  getDefaultFirstprivateAddressSpace());
-  cast<llvm::GlobalValue>(Addr)->setLinkage(Linkage);
-  CharUnits VarSize = CGM.getContext().getTypeSizeInChars(Ty);
-  CGM.addCompilerUsedGlobal(cast<llvm::GlobalValue>(Addr));
-  OffloadEntriesInfoManager.registerDeviceGlobalVarEntryInfo(
-      VarName, Addr, VarSize,
-      OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryTo, Linkage);
-  return Addr;
-}
-
 void CGOpenMPRuntime::registerTargetGlobalVariable(const VarDecl *VD,
                                                    llvm::Constant *Addr) {
   if (CGM.getLangOpts().OMPTargetTriples.empty() &&
       !CGM.getLangOpts().OpenMPIsDevice)
     return;
 
   // If we have host/nohost variables, they do not need to be registered.
   Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

public:
   /// the variables captured by reference in lambdas.
   void adjustTargetSpecificDataForLambdas(
       CodeGenFunction &CGF, const OMPExecutableDirective &D) const override;
 
   /// Perform check on requires decl to ensure that target architecture
   /// supports unified addressing
   void processRequiresDirective(const OMPRequiresDecl *D) override;
 
-  /// Returns default address space for the constant firstprivates, __constant__
-  /// address space by default.
-  unsigned getDefaultFirstprivateAddressSpace() const override;
-
   /// Checks if the variable has associated OMPAllocateDeclAttr attribute with
   /// the predefined allocator and translates it into the corresponding address
   /// space.
   bool hasAllocateAttributeForGlobalVar(const VarDecl *VD, LangAS &AS) override;
 
 private:
   /// Track the execution mode when codegening directives within a target
   /// region. The appropriate mode (SPMD/NON-SPMD) is set on entry to the

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

CGOpenMPRuntimeGPU::translateParameter(const FieldDecl *FD,
   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);
-    } else if (Attr->getCaptureKind() == OMPC_firstprivate &&
-               PointeeTy.isConstant(CGM.getContext())) {
-      PointeeTy = CGM.getContext().getAddrSpaceQualType(PointeeTy,
-                                                        LangAS::opencl_generic);
     }
   }
   ArgType = CGM.getContext().getPointerType(PointeeTy);
   QC.addRestrict();
   enum { NVPTX_local_addr = 5 };
   QC.addAddressSpace(getLangASFromTargetAS(NVPTX_local_addr));
   ArgType = QC.apply(CGM.getContext(), ArgType);
   if (isa<ImplicitParamDecl>(NativeParam))
for (const LambdaCapture &LC : RD->captures()) {
         VDAddr = CGF.EmitLoadOfReferenceLValue(VDAddr,
                                                VD->getType().getCanonicalType())
                      .getAddress(CGF);
       CGF.EmitStoreOfScalar(VDAddr.getPointer(), VarLVal);
     }
   }
 }
 
-unsigned CGOpenMPRuntimeGPU::getDefaultFirstprivateAddressSpace() const {
-  return CGM.getContext().getTargetAddressSpace(LangAS::cuda_constant);
-}
-
 bool CGOpenMPRuntimeGPU::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
                                                             LangAS &AS) {
   if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
     return false;
   const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
   switch(A->getAllocatorType()) {
   case OMPAllocateDeclAttr::OMPNullMemAlloc:
   case OMPAllocateDeclAttr::OMPDefaultMemAlloc:

clang/lib/CodeGen/CGStmtOpenMP.cpp

if (OriginalType->isArrayType()) {
     EmitIgnoredExpr(Copy);
   }
 }
 
 bool CodeGenFunction::EmitOMPFirstprivateClause(const OMPExecutableDirective &D,
                                                 OMPPrivateScope &PrivateScope) {
   if (!HaveInsertPoint())
     return false;
-  bool DeviceConstTarget =
-      getLangOpts().OpenMPIsDevice &&
-      isOpenMPTargetExecutionDirective(D.getDirectiveKind());
   bool FirstprivateIsLastprivate = false;
   llvm::DenseMap<const VarDecl *, OpenMPLastprivateModifier> Lastprivates;
   for (const auto *C : D.getClausesOfKind<OMPLastprivateClause>()) {
     for (const auto *D : C->varlists())
       Lastprivates.try_emplace(
           cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl(),
           C->getKind());
   }
for (const Expr *IInit : C->private_copies()) {
       if (!MustEmitFirstprivateCopy && !ThisFirstprivateIsLastprivate && FD &&
           !FD->getType()->isReferenceType() &&
           (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())) {
         EmittedAsFirstprivate.insert(OrigVD->getCanonicalDecl());
         ++IRef;
         ++InitsRef;
         continue;
       }
-      // Do not emit copy for firstprivate constant variables in target regions,
-      // captured by reference.
-      if (DeviceConstTarget && OrigVD->getType().isConstant(getContext()) &&
-          FD && FD->getType()->isReferenceType() &&
-          (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())) {
-        (void)CGM.getOpenMPRuntime().registerTargetFirstprivateCopy(*this,
-                                                                    OrigVD);
-        ++IRef;
-        ++InitsRef;
-        continue;
-      }
       FirstprivateIsLastprivate =
           FirstprivateIsLastprivate || ThisFirstprivateIsLastprivate;
       if (EmittedAsFirstprivate.insert(OrigVD->getCanonicalDecl()).second) {
         const auto *VDInit =
             cast<VarDecl>(cast<DeclRefExpr>(*InitsRef)->getDecl());
         bool IsRegistered;
         DeclRefExpr DRE(getContext(), const_cast<VarDecl *>(OrigVD),
                         /*RefersToEnclosingVariableOrCapture=*/FD != nullptr,

clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp

struct TT {
   ty Y;
 };
 
 // TCHECK-DAG:  [[TTII:%.+]] = type { i32, i32 }
 // TCHECK-DAG:  [[TTIC:%.+]] = type { i8, i8 }
 // TCHECK-DAG:  [[TT:%.+]] = type { i64, i8 }
 // TCHECK-DAG:  [[S1:%.+]] = type { double }
 
-// TCHECK: @__omp_offloading_firstprivate__{{.+}}_e_l30 = internal addrspace(4) global [[TTII]] zeroinitializer
-// TCHECK: @__omp_offloading_firstprivate__{{.+}}_ZTSK2TTIiiE_t_l143 = internal addrspace(4) global [[TTII]] zeroinitializer
-// TCHECK: @__omp_offloading_firstprivate__{{.+}}_ZTSK2TTIccE_t_l143 = internal addrspace(4) global [[TTIC]] zeroinitializer
 int foo(int n, double *ptr) {
   int a = 0;
   short aa = 0;
   float b[10];
   double c[5][10];
   TT<long long, char> d;
   const TT<int, int> e = {n, n};
 
 #pragma omp target firstprivate(a, e) map(tofrom \
                                           : b)
   {
     b[a] = a;
     b[a] += e.X;
   }
 
   // TCHECK:  define {{.*}}void @__omp_offloading_{{.+}}([10 x float] addrspace(1)* noalias [[B_IN:%.+]], i{{[0-9]+}} [[A_IN:%.+]], [[TTII]]* noalias [[E_IN:%.+]])
-  // TCHECK-NOT: alloca [[TTII]],
   // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
-  // TCHECK-NOT: alloca [[TTII]],
-  // TCHECK-NOT: alloca i{{[0-9]+}},
-  // TCHECK-64:  call void @llvm.dbg.declare(metadata [10 x float] addrspace(1)** %{{.+}}, metadata !{{[0-9]+}}, metadata !DIExpression())
+  // TCHECK: alloca [[TTII]],
+  // TCHECK: alloca i{{[0-9]+}},
   // TCHECK:  store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
   // TCHECK:  ret void
 
 #pragma omp target firstprivate(aa, b, c, d)
   {
     aa += 1;
     b[2] = 1.0;
     c[1][2] = 1.0;

clang/test/OpenMP/target_firstprivate_codegen.cpp

 // CHECK-DAG:  [[TT:%.+]] = type { i64, i8 }
 // CHECK-DAG:  [[TTII:%.+]] = type { i32, i32 }
 // CHECK-DAG:  [[S1:%.+]] = type { double }
 
 // TCHECK-DAG:  [[TT:%.+]] = type { i64, i8 }
 // TCHECK-DAG:  [[TTII:%.+]] = type { i32, i32 }
 // TCHECK-DAG:  [[S1:%.+]] = type { double }
 
-// CHECK-DAG:  [[FP_E:@__omp_offloading_firstprivate_.+_e_l79]] = internal global [[TTII]] zeroinitializer
 // CHECK-DAG:  [[SIZET:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ:32|64]] 4, i{{64|32}} {{8|4}}, i[[SZ:32|64]] 4]
 // CHECK-DAG:  [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 288, i64 49, i64 288]
 // CHECK-DAG:  [[MAPT2:@.+]] = private unnamed_addr constant [9 x i64] [i64 288, i64 161, i64 800, i64 161, i64 161, i64 800, i64 800, i64 161, i64 161]
 // CHECK-DAG:  [[SIZET3:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i{{32|64}} 0, i{{32|64}} 8]
-// CHECK-DAG:  [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 37]
+// CHECK-DAG:  [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 161]
 // CHECK-DAG:  [[MAPT4:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 800, i64 800, i64 161]
 // CHECK-DAG:  [[SIZET5:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG:  [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 288, i64 288, i64 161]
 // CHECK-DAG:  [[SIZET6:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 40]
 // CHECK-DAG:  [[MAPT6:@.+]] = private unnamed_addr constant [2 x i64] [i64 288, i64 161]
 
 // CHECK: define {{.*}}[[FOO:@.+]](
 int foo(int n, double *ptr) {
#pragma omp target firstprivate(a, p, ga)
   // CHECK:  [[N_ADDR:%.+]] = alloca i{{[0-9]+}},
   // CHECK:  [[PTR_ADDR:%.+]] = alloca double*,
   // CHECK:  [[A:%.+]] = alloca i{{[0-9]+}},
   // CHECK:  [[A2:%.+]] = alloca i{{[0-9]+}},
   // CHECK:  [[B:%.+]] = alloca [10 x float],
   // CHECK:  [[SSTACK:%.+]] = alloca i8*,
   // CHECK:  [[C:%.+]] = alloca [5 x [10 x double]],
   // CHECK:  [[D:%.+]] = alloca [[TT]],
+  // CHECK:  [[FP_E:%.+]] = alloca [[TTII]],
   // CHECK:  [[P:%.+]] = alloca i32*, align 64
   // CHECK:  [[ACAST:%.+]] = alloca i{{[0-9]+}},
   // CHECK:  [[BASE_PTR_ARR:%.+]] = alloca [3 x i8*],
   // CHECK:  [[PTR_ARR:%.+]] = alloca [3 x i8*],
   // CHECK:  [[A2CAST:%.+]] = alloca i{{[0-9]+}},
   // CHECK:  [[BASE_PTR_ARR2:%.+]] = alloca [9 x i8*],
   // CHECK:  [[PTR_ARR2:%.+]] = alloca [9 x i8*],
   // CHECK:  [[SIZET2:%.+]] = alloca [9 x i{{[0-9]+}}],

 
 #pragma omp target firstprivate(ptr, e)
   {
     ptr[0] = e.X;
     ptr[0]++;
   }
   // CHECK:  [[PTR_ADDR_REF:%.+]] = load double*, double** [[PTR_ADDR]],
 
-  // CHECK:  [[E_BC:%.+]] = bitcast [[TTII]]* [[E:%.+]] to i8*
-  // CHECK:  call void @llvm.memcpy.p0i8.p0i8.i{{64|32}}(i8* {{.*}} bitcast ([[TTII]]* [[FP_E]] to i8*), i8* {{.*}} [[E_BC]], i{{64|32}} 8, i1 false)
   // CHECK:  [[BASE_PTR_GEP3_0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
   // CHECK:  [[BCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTR_GEP3_0]] to double**
   // CHECK:  store double* [[PTR_ADDR_REF]], double** [[BCAST_TOPTR]],
   // CHECK:  [[PTR_GEP3_0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
   // CHECK:  [[BCAST_TOPTR:%.+]] = bitcast i8** [[PTR_GEP3_0]] to double**
   // CHECK:  store double* [[PTR_ADDR_REF]], double** [[BCAST_TOPTR]],
   // CHECK:  [[BASE_PTR_GEP3_1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
   // CHECK:  [[BCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTR_GEP3_1]] to [[TTII]]**
   // CHECK:  store [[TTII]]* [[FP_E]], [[TTII]]** [[BCAST_TOPTR]],
   // CHECK:  [[PTR_GEP3_1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
   // CHECK:  [[BCAST_TOPTR:%.+]] = bitcast i8** [[PTR_GEP3_1]] to [[TTII]]**
   // CHECK:  store [[TTII]]* [[FP_E]], [[TTII]]** [[BCAST_TOPTR]],
 
   // CHECK:  [[BASE_PTR_GEP_ARG3:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
   // CHECK:  [[PTR_GEP_ARG3:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
   // CHECK: {{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, {{.+}}, i32 2, i8** [[BASE_PTR_GEP_ARG3]], i8** [[PTR_GEP_ARG3]], i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET3]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT3]], i32 0, i32 0), i8** null, i8** null)
 
   // TCHECK:  define weak void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]], [[TTII]]* nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[E:%.+]])
-  // TCHECK-NOT: alloca [[TTII]],
   // TCHECK:  [[PTR_ADDR:%.+]] = alloca double*,
-  // TCHECK-NOT: alloca [[TTII]],
+  // TCHECK: alloca [[TTII]],
   // TCHECK-NOT: alloca double*,
   // TCHECK:  store double* [[PTR_IN]], double** [[PTR_ADDR]],
   // TCHECK-NOT: store double* %
 
   return a;
 }
 
 template <typename tx>