void linalg::convertLinalg3ToLLVM(Module module) {
// Remove affine constructs.
- for (auto func : module) {
+ for (auto func : module.getOps<FuncOp>()) {
auto rr = lowerAffineConstructs(func);
(void)rr;
assert(succeeded(rr) && "affine loop lowering failed");
// Delete any generic function left
// FIXME: we may want this as a separate pass.
- for (mlir::Function function : llvm::make_early_inc_range(module)) {
+ for (mlir::Function function :
+ llvm::make_early_inc_range(module.getOps<mlir::Function>())) {
if (auto genericAttr =
function.getAttrOfType<mlir::BoolAttr>("toy.generic")) {
if (genericAttr.getValue())
// affine dialect: they already include conversion to the LLVM dialect.
// First patch calls type to return memref instead of ToyArray
- for (auto function : getModule()) {
+ for (auto function : getModule().getOps<FuncOp>()) {
function.walk([&](Operation *op) {
auto callOp = dyn_cast<CallOp>(op);
if (!callOp)
});
}
- for (auto function : getModule()) {
+ for (auto function : getModule().getOps<FuncOp>()) {
function.walk([&](Operation *op) {
// Turns toy.alloc into sequence of alloc/dealloc (later malloc/free).
if (auto allocOp = dyn_cast<toy::AllocOp>(op)) {
}
// Lower Linalg to affine
- for (auto function : getModule())
+ for (auto function : getModule().getOps<FuncOp>())
linalg::lowerToLoops(function);
getModule().dump();
// Delete any generic function left
// FIXME: we may want this as a separate pass.
- for (mlir::Function function : llvm::make_early_inc_range(module)) {
+ for (mlir::Function function :
+ llvm::make_early_inc_range(module.getOps<mlir::Function>())) {
if (auto genericAttr =
function.getAttrOfType<mlir::BoolAttr>("toy.generic")) {
if (genericAttr.getValue())
// Body Management.
//===--------------------------------------------------------------------===//
- // Iterate over the functions within the module.
- using iterator = Block::op_iterator<FuncOp>;
-
- // Iteration over the functions in the module.
- iterator begin() { return getBody()->op_begin<FuncOp>(); }
- iterator end() { return getBody()->op_end<FuncOp>(); }
- Function front() { return *begin(); }
- Function back() { return *std::prev(end()); }
-
- /// This is the list of functions in the module.
- llvm::iterator_range<iterator> getFunctions() {
- return getBody()->getOps<FuncOp>();
+ /// Iteration over the operations in the module.
+ using iterator = Block::iterator;
+
+ iterator begin() { return getBody()->begin(); }
+ iterator end() { return getBody()->end(); }
+ Operation &front() { return *begin(); }
+
+ /// This returns a range of operations of the given type 'T' held within the
+ /// module.
+ template <typename T> llvm::iterator_range<Block::op_iterator<T>> getOps() {
+ return getBody()->getOps<T>();
}
/// Insert the operation into the back of the body, before the terminator.
insert(Block::iterator(getBody()->getTerminator()), op);
}
- /// Inser the operation at the given insertion point. Note: The operation is
+ /// Insert the operation at the given insertion point. Note: The operation is
/// never inserted after the terminator, even if the insertion point is end().
void insert(Operation *insertPt, Operation *op) {
insert(Block::iterator(insertPt), op);
/// name exists. Function names never include the @ on them. Note: This
/// performs a linear scan of held symbols.
Function getNamedFunction(Identifier name) {
- auto it = llvm::find_if(getFunctions(),
+ auto it = llvm::find_if(getOps<FuncOp>(),
[name](FuncOp fn) { return fn.getName() == name; });
return it == end() ? nullptr : *it;
}
opCount.clear();
// Compute the operation statistics for each function in the module.
- for (auto fn : getModule())
- fn.walk([&](Operation *op) { ++opCount[op->getName().getStringRef()]; });
+ for (auto &op : getModule())
+ op.walk([&](Operation *op) { ++opCount[op->getName().getStringRef()]; });
printSummary();
}
LLVMInitializeNVPTXTargetMC();
LLVMInitializeNVPTXAsmPrinter();
- for (auto function : getModule()) {
+ for (auto function : getModule().getOps<FuncOp>()) {
if (!gpu::GPUDialect::isKernel(function) || function.isExternal()) {
continue;
}
// Cache the used LLVM types.
initializeCachedTypes();
- for (auto func : getModule()) {
+ for (auto func : getModule().getOps<FuncOp>()) {
func.walk<mlir::gpu::LaunchFuncOp>(
[this](mlir::gpu::LaunchFuncOp op) { translateGpuLaunchCalls(op); });
}
auto module = getModule();
Builder builder(&getContext());
- auto functions = module.getFunctions();
+ auto functions = module.getOps<FuncOp>();
for (auto it = functions.begin(); it != functions.end();) {
// Move iterator to after the current function so that potential insertion
// of the accessor is after the kernel with cubin iself.
}
void mlir::LLVM::ensureDistinctSuccessors(Module m) {
- for (auto f : m) {
+ for (auto f : m.getOps<FuncOp>()) {
for (auto &bb : f.getBlocks()) {
::ensureDistinctSuccessors(bb);
}
public:
void runOnModule() override {
ModuleManager moduleManager(getModule());
- for (auto func : getModule()) {
+ for (auto func : getModule().getOps<FuncOp>()) {
func.walk<mlir::gpu::LaunchOp>([&](mlir::gpu::LaunchOp op) {
Function outlinedFunc = outlineKernelFunc(op);
moduleManager.insert(outlinedFunc);
// Check that all functions are uniquely named.
llvm::StringMap<Location> nameToOrigLoc;
- for (auto fn : getFunctions()) {
+ for (auto fn : getOps<FuncOp>()) {
auto it = nameToOrigLoc.try_emplace(fn.getName(), fn.getLoc());
if (!it.second)
return fn.emitError()
/// Build a symbol table with the symbols within the given module.
SymbolTable::SymbolTable(ModuleOp module) : context(module.getContext()) {
- for (auto func : module) {
+ for (auto func : module.getOps<FuncOp>()) {
auto inserted = symbolTable.insert({func.getName(), func});
(void)inserted;
assert(inserted.second &&
void LowerLinalgToLLVMPass::runOnModule() {
auto module = getModule();
- for (auto f : module.getFunctions()) {
+ for (auto f : module.getOps<FuncOp>()) {
lowerLinalgSubViewOps(f);
lowerLinalgForToCFG(f);
if (failed(lowerAffineConstructs(f)))
/// module.
void ModuleToFunctionPassAdaptor::runOnModule() {
ModuleAnalysisManager &mam = getAnalysisManager();
- for (auto func : getModule()) {
+ for (auto func : getModule().getOps<FuncOp>()) {
// Skip external functions.
if (func.isExternal())
continue;
// This ensures that an analysis manager exists for each function, as well as
// providing a queue of functions to execute over.
std::vector<std::pair<Function, FunctionAnalysisManager>> funcAMPairs;
- for (auto func : getModule())
+ for (auto func : getModule().getOps<FuncOp>())
if (!func.isExternal())
funcAMPairs.emplace_back(func, mam.slice(func));
void InferQuantizedTypesPass::runWithConfig(SolverContext &solverContext,
const TargetConfiguration &config) {
CAGSlice cag(solverContext);
- for (auto f : getModule()) {
+ for (auto f : getModule().getOps<FuncOp>()) {
f.walk([&cag, &config](Operation *op) { config.handleOp(op, cag); });
}
config.finalizeAnchors(cag);
// wrapping the SPIR-V ModuleOp inside a MLIR module. This should be changed
// to take in the SPIR-V ModuleOp directly after module and function are
// migrated to be general ops.
- for (auto fn : module) {
+ for (auto fn : module.getOps<FuncOp>()) {
fn.walk<spirv::ModuleOp>([&](spirv::ModuleOp spirvModule) {
if (done) {
spirvModule.emitError("found more than one 'spv.module' op");
// Insert the nvvm.annotations kernel so that the NVVM backend recognizes the
// function as a kernel.
- for (Function func : m) {
+ for (Function func : m.getOps<FuncOp>()) {
if (!func.getAttrOfType<UnitAttr>(gpu::GPUDialect::getKernelFuncAttrName()))
continue;
bool ModuleTranslation::convertFunctions() {
// Declare all functions first because there may be function calls that form a
// call graph with cycles.
- for (Function function : mlirModule) {
+ for (Function function : mlirModule.getOps<FuncOp>()) {
mlir::BoolAttr isVarArgsAttr =
function.getAttrOfType<BoolAttr>("std.varargs");
bool isVarArgs = isVarArgsAttr && isVarArgsAttr.getValue();
}
// Convert functions.
- for (Function function : mlirModule) {
+ for (Function function : mlirModule.getOps<FuncOp>()) {
// Ignore external functions.
if (function.isExternal())
continue;
mlir::applyConversionPatterns(Module module, ConversionTarget &target,
TypeConverter &converter,
OwningRewritePatternList &&patterns) {
- SmallVector<Function, 32> allFunctions(module.getFunctions());
+ SmallVector<Function, 32> allFunctions(module.getOps<FuncOp>());
return applyConversionPatterns(allFunctions, target, converter,
std::move(patterns));
}
projects / platform / upstream / llvm.git / commitdiff