Merge branch 'llvm:main' into main

Author: code-pankaj

clang/include/clang/AST/VTableBuilder.h

@@ -246,12 +246,12 @@ class VTableLayout {
  // point for a given vtable index.
  typedef llvm::SmallVector<unsigned, 4> AddressPointsIndexMapTy;
 
+ using VTableIndicesTy = llvm::SmallVector<std::size_t>;
+
 private:
- // Stores the component indices of the first component of each virtual table in
- // the virtual table group. To save a little memory in the common case where
- // the vtable group contains a single vtable, an empty vector here represents
- // the vector {0}.
- OwningArrayRef<size_t> VTableIndices;
+ // Stores the component indices of the first component of each virtual table
+ // in the virtual table group.
+ VTableIndicesTy VTableIndices;
 
  OwningArrayRef<VTableComponent> VTableComponents;
 
@@ -265,7 +265,8 @@ class VTableLayout {
  AddressPointsIndexMapTy AddressPointIndices;
 
 public:
- VTableLayout(ArrayRef<size_t> VTableIndices,
+ // Requires `VTableIndices.front() == 0`
+ VTableLayout(VTableIndicesTy VTableIndices,
  ArrayRef<VTableComponent> VTableComponents,
  ArrayRef<VTableThunkTy> VTableThunks,
  const AddressPointsMapTy &AddressPoints);
@@ -292,26 +293,11 @@ class VTableLayout {
  return AddressPointIndices;
  }
 
- size_t getNumVTables() const {
- if (VTableIndices.empty())
- return 1;
- return VTableIndices.size();
- }
+ size_t getNumVTables() const { return VTableIndices.size(); }
 
- size_t getVTableOffset(size_t i) const {
- if (VTableIndices.empty()) {
- assert(i == 0);
- return 0;
- }
- return VTableIndices[i];
- }
+ size_t getVTableOffset(size_t i) const { return VTableIndices[i]; }
 
  size_t getVTableSize(size_t i) const {
- if (VTableIndices.empty()) {
- assert(i == 0);
- return vtable_components().size();
- }
-
  size_t thisIndex = VTableIndices[i];
  size_t nextIndex = (i + 1 == VTableIndices.size())
  ? vtable_components().size()

clang/lib/AST/VTableBuilder.cpp

@@ -999,7 +999,7 @@ class ItaniumVTableBuilder {
 public:
  /// Component indices of the first component of each of the vtables in the
  /// vtable group.
- SmallVector<size_t, 4> VTableIndices;
+ VTableLayout::VTableIndicesTy VTableIndices;
 
  ItaniumVTableBuilder(ItaniumVTableContext &VTables,
  const CXXRecordDecl *MostDerivedClass,
@@ -2306,18 +2306,19 @@ MakeAddressPointIndices(const VTableLayout::AddressPointsMapTy &addressPoints,
  return indexMap;
 }
 
-VTableLayout::VTableLayout(ArrayRef<size_t> VTableIndices,
+VTableLayout::VTableLayout(VTableIndicesTy VTableIndices,
  ArrayRef<VTableComponent> VTableComponents,
  ArrayRef<VTableThunkTy> VTableThunks,
  const AddressPointsMapTy &AddressPoints)
- : VTableComponents(VTableComponents), VTableThunks(VTableThunks),
- AddressPoints(AddressPoints), AddressPointIndices(MakeAddressPointIndices(
- AddressPoints, VTableIndices.size())) {
- if (VTableIndices.size() <= 1)
- assert(VTableIndices.size() == 1 && VTableIndices[0] == 0);
- else
- this->VTableIndices = OwningArrayRef<size_t>(VTableIndices);
-
+ : VTableIndices(std::move(VTableIndices)),
+ VTableComponents(VTableComponents), VTableThunks(VTableThunks),
+ AddressPoints(AddressPoints),
+ AddressPointIndices(
+ MakeAddressPointIndices(AddressPoints, this->VTableIndices.size())) {
+ assert(!this->VTableIndices.empty() &&
+ "VTableLayout requires at least one index.");
+ assert(this->VTableIndices.front() == 0 &&
+ "VTableLayout requires the first index is 0.");
  llvm::sort(this->VTableThunks, [](const VTableLayout::VTableThunkTy &LHS,
  const VTableLayout::VTableThunkTy &RHS) {
  assert((LHS.first != RHS.first || LHS.second == RHS.second) &&
@@ -3730,8 +3731,8 @@ void MicrosoftVTableContext::computeVTableRelatedInformation(
  SmallVector<VTableLayout::VTableThunkTy, 1> VTableThunks(
  Builder.vtable_thunks_begin(), Builder.vtable_thunks_end());
  VFTableLayouts[id] = std::make_unique<VTableLayout>(
- ArrayRef<size_t>{0}, Builder.vtable_components(), VTableThunks,
- EmptyAddressPointsMap);
+ VTableLayout::VTableIndicesTy{0}, Builder.vtable_components(),
+ VTableThunks, EmptyAddressPointsMap);
  Thunks.insert(Builder.thunks_begin(), Builder.thunks_end());
 
  const ASTRecordLayout &Layout = Context.getASTRecordLayout(RD);

clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp

@@ -306,13 +306,15 @@ class OpenACCRoutineClauseEmitter final
  CIRGenModule &cgm;
  CIRGen::CIRGenBuilderTy &builder;
  mlir::acc::RoutineOp routineOp;
+ const clang::FunctionDecl *funcDecl;
  llvm::SmallVector<mlir::acc::DeviceType> lastDeviceTypeValues;
 
 public:
  OpenACCRoutineClauseEmitter(CIRGenModule &cgm,
  CIRGen::CIRGenBuilderTy &builder,
- mlir::acc::RoutineOp routineOp)
- : cgm(cgm), builder(builder), routineOp(routineOp) {}
+ mlir::acc::RoutineOp routineOp,
+ const clang::FunctionDecl *funcDecl)
+ : cgm(cgm), builder(builder), routineOp(routineOp), funcDecl(funcDecl) {}
 
  void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
  this->VisitClauseList(clauses);
@@ -372,8 +374,12 @@ class OpenACCRoutineClauseEmitter final
  value);
  } else {
  assert(clause.isIdentifierArgument());
- cgm.errorNYI(clause.getSourceRange(),
- "Bind with an identifier argument is not yet supported");
+ std::string bindName = cgm.getOpenACCBindMangledName(
+ clause.getIdentifierArgument(), funcDecl);
+
+ routineOp.addBindIDName(
+ builder.getContext(), lastDeviceTypeValues,
+ mlir::SymbolRefAttr::get(builder.getContext(), bindName));
  }
  }
 };
@@ -416,6 +422,6 @@ void CIRGenModule::emitOpenACCRoutineDecl(
  mlir::acc::getRoutineInfoAttrName(),
  mlir::acc::RoutineInfoAttr::get(func.getContext(), funcRoutines));
 
- OpenACCRoutineClauseEmitter emitter{*this, builder, routineOp};
+ OpenACCRoutineClauseEmitter emitter{*this, builder, routineOp, funcDecl};
  emitter.emitClauses(clauses);
 }

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -1721,6 +1721,71 @@ static std::string getMangledNameImpl(CIRGenModule &cgm, GlobalDecl gd,
  return std::string(out.str());
 }
 
+static FunctionDecl *
+createOpenACCBindTempFunction(ASTContext &ctx, const IdentifierInfo *bindName,
+ const FunctionDecl *protoFunc) {
+ // If this is a C no-prototype function, we can take the 'easy' way out and
+ // just create a function with no arguments/functions, etc.
+ if (!protoFunc->hasPrototype())
+ return FunctionDecl::Create(
+ ctx, /*DC=*/ctx.getTranslationUnitDecl(),
+ /*StartLoc=*/SourceLocation{}, /*NLoc=*/SourceLocation{}, bindName,
+ protoFunc->getType(), /*TInfo=*/nullptr, StorageClass::SC_None);
+
+ QualType funcTy = protoFunc->getType();
+ auto *fpt = cast<FunctionProtoType>(protoFunc->getType());
+
+ // If this is a member function, add an explicit 'this' to the function type.
+ if (auto *methodDecl = dyn_cast<CXXMethodDecl>(protoFunc);
+ methodDecl && methodDecl->isImplicitObjectMemberFunction()) {
+ llvm::SmallVector<QualType> paramTypes{fpt->getParamTypes()};
+ paramTypes.insert(paramTypes.begin(), methodDecl->getThisType());
+
+ funcTy = ctx.getFunctionType(fpt->getReturnType(), paramTypes,
+ fpt->getExtProtoInfo());
+ fpt = cast<FunctionProtoType>(funcTy);
+ }
+
+ auto *tempFunc =
+ FunctionDecl::Create(ctx, /*DC=*/ctx.getTranslationUnitDecl(),
+ /*StartLoc=*/SourceLocation{},
+ /*NLoc=*/SourceLocation{}, bindName, funcTy,
+ /*TInfo=*/nullptr, StorageClass::SC_None);
+
+ SmallVector<ParmVarDecl *> params;
+ params.reserve(fpt->getNumParams());
+
+ // Add all of the parameters.
+ for (unsigned i = 0, e = fpt->getNumParams(); i != e; ++i) {
+ ParmVarDecl *parm = ParmVarDecl::Create(
+ ctx, tempFunc, /*StartLoc=*/SourceLocation{},
+ /*IdLoc=*/SourceLocation{},
+ /*Id=*/nullptr, fpt->getParamType(i), /*TInfo=*/nullptr,
+ StorageClass::SC_None, /*DefArg=*/nullptr);
+ parm->setScopeInfo(0, i);
+ params.push_back(parm);
+ }
+
+ tempFunc->setParams(params);
+
+ return tempFunc;
+}
+
+std::string
+CIRGenModule::getOpenACCBindMangledName(const IdentifierInfo *bindName,
+ const FunctionDecl *attachedFunction) {
+ FunctionDecl *tempFunc = createOpenACCBindTempFunction(
+ getASTContext(), bindName, attachedFunction);
+
+ std::string ret = getMangledNameImpl(*this, GlobalDecl(tempFunc), tempFunc);
+
+ // This does nothing (it is a do-nothing function), since this is a
+ // slab-allocator, but leave a call in to immediately destroy this in case we
+ // ever come up with a way of getting allocations back.
+ getASTContext().Deallocate(tempFunc);
+ return ret;
+}
+
 StringRef CIRGenModule::getMangledName(GlobalDecl gd) {
  GlobalDecl canonicalGd = gd.getCanonicalDecl();
 

clang/lib/CIR/CodeGen/CIRGenModule.h

@@ -507,6 +507,15 @@ class CIRGenModule : public CIRGenTypeCache {
  mlir::Value emitMemberPointerConstant(const UnaryOperator *e);
 
  llvm::StringRef getMangledName(clang::GlobalDecl gd);
+ // This function is to support the OpenACC 'bind' clause, which names an
+ // alternate name for the function to be called by. This function mangles
+ // `attachedFunction` as-if its name was actually `bindName` (that is, with
+ // the same signature). It has some additional complications, as the 'bind'
+ // target is always going to be a global function, so member functions need an
+ // explicit instead of implicit 'this' parameter, and thus gets mangled
+ // differently.
+ std::string getOpenACCBindMangledName(const IdentifierInfo *bindName,
+ const FunctionDecl *attachedFunction);
 
  void emitTentativeDefinition(const VarDecl *d);
 

clang/test/CIR/CodeGenOpenACC/routine-bind.c

@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+// FIXME: We should run this against Windows mangling as well at one point.
 
 #pragma acc routine seq bind("BIND1")
 void Func1(){}
@@ -18,6 +19,28 @@ void Func5(){}
 void Func6(){}
 #pragma acc routine(Func6) seq device_type(radeon) bind("BIND6_R") device_type(multicore, host) bind("BIND6_M")
 
+#pragma acc routine seq bind(BIND7)
+void Func7(int i){}
+
+void Func8(float f){}
+#pragma acc routine(Func8) seq bind(BIND8)
+
+#pragma acc routine seq device_type(nvidia) bind(BIND9)
+void Func9(int i, float f, short s){}
+
+struct S{};
+struct U{};
+struct V{};
+
+void Func10(struct S s){}
+#pragma acc routine(Func10) seq device_type(radeon) bind(BIND10)
+
+#pragma acc routine seq device_type(nvidia, host) bind(BIND11_NVH) device_type(multicore) bind(BIND11_MC)
+void Func11(struct U* u, struct V v, int i){}
+
+int Func12(struct U u, struct V v, int i){ return 0; }
+#pragma acc routine(Func12) seq device_type(radeon) bind(BIND12_R) device_type(multicore, host) bind(BIND12_MCH)
+
 // CHECK: cir.func{{.*}} @[[F1_NAME:.*Func1[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>}
 // CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) bind("BIND1") seq
 //
@@ -33,7 +56,25 @@ void Func6(){}
 //
 // CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>}
 //
+// CHECK: cir.func{{.*}} @[[F7_NAME:.*Func7[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F7_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F7_R_NAME]] func(@[[F7_NAME]]) bind(@BIND7) seq
+//
+// CHECK: cir.func{{.*}} @[[F8_NAME:.*Func8[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F8_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[F9_NAME:.*Func9[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F9_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F9_R_NAME]] func(@[[F9_NAME]]) bind(@BIND9 [#acc.device_type<nvidia>]) seq
+//
+// CHECK: cir.func{{.*}} @[[F10_NAME:.*Func10[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F10_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[F11_NAME:.*Func11[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F11_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F11_R_NAME]] func(@[[F11_NAME]]) bind(@BIND11_NVH [#acc.device_type<nvidia>], @BIND11_NVH [#acc.device_type<host>], @BIND11_MC [#acc.device_type<multicore>])
+//
+// CHECK: cir.func{{.*}} @[[F12_NAME:.*Func12[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F12_R_NAME:.*]]]>}
+//
 // CHECK: acc.routine @[[F2_R_NAME]] func(@[[F2_NAME]]) bind("BIND2") seq
 // CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) bind("BIND4" [#acc.device_type<radeon>]) seq
 // CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) bind("BIND6_R" [#acc.device_type<radeon>], "BIND6_M" [#acc.device_type<multicore>], "BIND6_M" [#acc.device_type<host>]) seq
 
+// CHECK: acc.routine @[[F8_R_NAME]] func(@[[F8_NAME]]) bind(@BIND8) seq
+// CHECK: acc.routine @[[F10_R_NAME]] func(@[[F10_NAME]]) bind(@BIND10 [#acc.device_type<radeon>]) seq
+// CHECK: acc.routine @[[F12_R_NAME]] func(@[[F12_NAME]]) bind(@BIND12_R [#acc.device_type<radeon>], @BIND12_MCH [#acc.device_type<multicore>], @BIND12_MCH [#acc.device_type<host>]) seq