Add a new interface allowing to set a default dialect to be used for printing/parsing regions

[mehdi_amini]

Summary

Currently the builtin dialect is the default namespace used for parsing
and printing. As such module and func don't need to be prefixed.
In the case of some dialects that defines new regions for their own
purpose (like SpirV modules for example), it can be beneficial to
change the default dialect in order to improve readability.

[mehdi_amini]

Fix bufferization of the stream and add doc

[mehdi_amini]

Don't use the default dialect to strip prefix for generic operations
(also update tests)

[ftynse]

Won't this also drop the occurrences of the string that looks like a prefix from other places than operation name? Everything goes through the same stream AFAICS, so something like my.op() attributes { my.attribute = "this.may.use.the.word.builtin.because", my.builtin.enabled = 1} will get affected where it shouldn't.

mlir/include/mlir/IR/OpImplementation.h

969–971 ↗

(On Diff #363386)

mlir/lib/IR/AsmPrinter.cpp

2322–2326

2343–2346

I can't match this comment with the code below.

[mehdi_amini]

In D107236#2919597, @ftynse wrote:

Won't this also drop the occurrences of the string that looks like a prefix from other places than operation name? Everything goes through the same stream AFAICS, so something like my.op() attributes { my.attribute = "this.may.use.the.word.builtin.because", my.builtin.enabled = 1} will get affected where it shouldn't.

I don't think so because we only filter the beginning of the stream and nothing else. The filter is immediately clear whether it matches or not the beginning of the stream. The filter is also set immediately before printing an operation. But I'll extend the test to show this!

[mehdi_amini]

Address comments

mlir/lib/IR/AsmPrinter.cpp

2343–2346

Ah, this is some explanation on the correctness of the code, I'll try to reformulate.

[mehdi_amini]

Fix getDefaultDialectFromEnclosingOps

[mehdi_amini]

mlir/include/mlir/IR/OpImplementation.h

969–971 ↗

(On Diff #363386)

Actually this edit was incorrect: if would skip the check on the current op

[mehdi_amini]

(some more minor fixes)

[rriddle]

Nice!

mlir/include/mlir/IR/OpAsmInterface.td

68

Can we instead just move this in the other interface? I don't think we need a special interface for this, especially given that OpAsmOpInterface is described as being a general parser/printer hook.

mlir/include/mlir/IR/OpImplementation.h

968 ↗

(On Diff #363584)

Can we move this to AsmPrinter.cpp? I don't see any uses outside of there, or is there some reason to have this publicly exposable?

971 ↗

(On Diff #363584)

Do we want it to be inherited? I would almost prefer that the default only apply to the immediate regions. It seems less surprising that way, and more contextualized. For example assuming we inherit default names, if I split out a nested region op from a reduced test case, I might need to re-add operation namespaces that were filtered out from the parent (as the parent may now be different).

mlir/lib/IR/AsmPrinter.cpp

2310

We discussed a little offline, but I think it would be better to just remove the constraint that custom op printers print the operation name at the beginning. I don't see a real reason for them to do that, and it would also match the behavior of the parser (which doesn't parse the operation name, for likely obvious reasons).

2433

Do we need this to be std::string? I would assume that we could enforce StringRef here (getDefaultDialect also returns StringRef).

mlir/lib/Parser/ParserState.h

91

Same comment here, can we use StringRef instead of std::string?

[rriddle]

mlir/include/mlir/IR/OpAsmInterface.td

83

Do we want to allow multiple default namespaces? Some projects might have several that all kind of co-mingle with each other.

mlir/include/mlir/IR/OpImplementation.h

973 ↗

(On Diff #363584)

Do we need to default to builtin? Can we instead default to nothing?

mlir/lib/IR/AsmPrinter.cpp

2433

I think I missed it, but how is this parser stack different from the one in the ParserState?

[mehdi_amini]

Address River's comment

[mehdi_amini]

mlir/include/mlir/IR/OpAsmInterface.td

83

I can't imagine what it'd look like right now: we'd have to guarantee these dialects don't have any pair of op with the same name to avoid conflicts.
The main extension I could think of would be to take an integer as an input and return a different one per region attached to an op

mlir/include/mlir/IR/OpImplementation.h

971 ↗

(On Diff #363584)

Yeah I hesitated on this, we can start with a shallow model conservatively

973 ↗

(On Diff #363584)

That mean that we wouldn't be able to parse a top-level module right now though, do you want to change this now?

(I'd be fine with this but I would land this in two revision though)

mlir/lib/IR/AsmPrinter.cpp

2310

Yeah I looked into it, one wrinkle though is that something like Dialect::printOperation can fail and we fallback to the generic operation printer.
I'm not sure how to handle this?

2433

This is a printer stack, and the other is the parser one? They are equivalent otherwise.
Maybe I missed what you were asking?

[rriddle]

mlir/include/mlir/IR/OpImplementation.h

973 ↗

(On Diff #363584)

I think it's reasonable/fine to require the "builtin" namespace on the top-level module. I'm also okay with it being separate from this commit.

mlir/lib/IR/AsmPrinter.cpp

2310

Hmmm. Can we change Dialect::printOperation to act more like getParseOperationHook? That way we would know if it would succeed before calling it.

2433

Right right, missed that one was printer/one was parser.

mlir/lib/Parser/Parser.cpp

1846

Would it be better to keep StringRef here, but have another variable that is OperationName? That would remove the need to build a std::string, and the OperationName var could be passed in later when building the OperationState.

1869–1875

Shouldn't this one be .str()?

[bondhugula]

mlir/include/mlir/IR/OpAsmInterface.td

56

under a given operation -> under this operation

(cf. doc of getAsmResultNames above to align.)

56

dialect -> dialect prefix

57

Nit: Punctuation broken here.

58–59

if it was implementing this interface to return ... -> if this method returned spv.

^^

"it" is dangling here.

mlir/lib/IR/AsmPrinter.cpp

2316–2317

Nit: reflow to use width.

2361

Camel case?

2429

get -> gets / is

[bondhugula]

purpose (like SpirV modules for example), it can be beneficial to
change the default dialect in order to improve readability.

Does this really improve readability? Doesn't it create ambiguity for a reader as to which dialect's operation is being referring to (when you have the same names)? I might have misunderstood it but that's what the commit summary conveys to me.

[mehdi_amini]

Does this really improve readability? Doesn't it create ambiguity for a reader as to which dialect's operation is being referring to (when you have the same names)? I might have misunderstood it but that's what the commit summary conveys to me.

Yes you’re right. I think this is about how it gets used though.
I’d claim that it is less intrusive than the custom assembly that we allow (we actually had internal teams complaining about it: affine, linalg, scf, … all have different custom conventions).

If you remember the original affine functions, I don’t think operations were prefixed: it was later during generalization that we added the prefix. The prefix wasn’t for readability but for necessary disambiguating in the parser.

The idea here is that you can see it like an extension of custom assembly: instead of just controlling the immediate operation you can impact the immediately nested regions in a limited way.

My main motivation is the TensorFlow graph dialect right now: when I write tfg.graph { I’d like to avoid prefixing every single operation with tfg. inside the graph itself.
This isn’t a region that allows anything else anyway, but even if it did, the few outlier operation would be prefixed.

[mehdi_amini]

Address comments

[mehdi_amini]

mlir/include/mlir/IR/OpAsmInterface.td

57

I can't spot what you refer to here?

mlir/lib/IR/AsmPrinter.cpp

2310

(I'll look into this later)

2361

I am overriding the parent class from LLVM, can't do?

mlir/lib/Parser/Parser.cpp

1846

I tried but it's convoluted: OperationName does not have a default constructor.

[jpienaar]

In D107236#2924729, @mehdi_amini wrote:

Does this really improve readability? Doesn't it create ambiguity for a reader as to which dialect's operation is being referring to (when you have the same names)? I might have misunderstood it but that's what the commit summary conveys to me.

If you remember the original affine functions, I don’t think operations were prefixed: it was later during generalization that we added the prefix. The prefix wasn’t for readability but for necessary disambiguating in the parser.

It did also disambiguate for the user though I'd say. Less of an issue with smaller number of dialects and all kept in head (we've also never had two dialects with i32 type for example, but if one saw i32 type one would not consider if from dialect or not)

The idea here is that you can see it like an extension of custom assembly: instead of just controlling the immediate operation you can impact the immediately nested regions in a limited way.

I think that is the question. This is like having a using namespace statement implicitly and attached to a region (well I think all regions of dialect specified ops?). So when looking at a dump a reader has to know which disambiguating rules to follow where implicitly (there is no visual cue). I'm assuming when an error is emitted we don't strip prefix (so that error is unambiguous without additional context).

And then Rivers question also come in to play: what if I have 3 dialects that I use equally in a given op's "first layer of ops" (as even TFG I believe allows nesting and ops inside the nested region may not be TFG), now do I have to decide based on preferred one.

I do agree that folks have complained about custom syntax being inconsistent, but that doesn't seem like a good argument for why this is good.

[mehdi_amini]

I do agree that folks have complained about custom syntax being inconsistent, but that doesn't seem like a good argument for why this is good.

Absolutely: the argument I made is that it is consistent with how we approached it in MLIR so far ; when the generic printer isn’t used you get an output that is readable optimally for folks familiar with a given dialect. And also that we provide this power to dialect authors, and trust them to use it responsibly.

(as even TFG I believe allows nesting and ops inside the nested region may not be TFG),

Nesting further under TFG shouldn’t be an issue, it would be a different region which may or may not use the same default (or none).

now do I have to decide based on preferred one.

Yes, or none: if three dialects are equally used in a region, why strip any of them? The fact that we can strip does not mean we should!
Other than Spirv that already restricts that only spirv ops are present, I don’t see any other place in-tree where we would use this feature right now.

[bondhugula]

In D107236#2924729, @mehdi_amini wrote:

Does this really improve readability? Doesn't it create ambiguity for a reader as to which dialect's operation is being referring to (when you have the same names)? I might have misunderstood it but that's what the commit summary conveys to me.

Yes you’re right. I think this is about how it gets used though.
I’d claim that it is less intrusive than the custom assembly that we allow (we actually had internal teams complaining about it: affine, linalg, scf, … all have different custom conventions).

If you remember the original affine functions, I don’t think operations were prefixed: it was later during generalization that we added the prefix. The prefix wasn’t for readability but for necessary disambiguating in the parser.

The idea here is that you can see it like an extension of custom assembly: instead of just controlling the immediate operation you can impact the immediately nested regions in a limited way.

My main motivation is the TensorFlow graph dialect right now: when I write tfg.graph { I’d like to avoid prefixing every single operation with tfg. inside the graph itself.
This isn’t a region that allows anything else anyway, but even if it did, the few outlier operation would be prefixed.

Then, the ambiguity that this adds due to the prefix being dropped based on the context appears really undesirable to me at first sight. It's more readable with the prefix for me. Of course, you are providing an option for a dialect to elide things, but this whole context sensitive printing (esp. for large regions) in the face of collisions can be quite painful. This also adds a layer of non-uniformity and surprises for people working across dialects and when folks encounter a new dialect. MLIR being an IR and not a user-facing programming language, I'm not sure you gain much by dropping a prefix like tfg to your ops.

Could you discuss this change on discourse? (Sorry if you already did and I missed it.) I feel we should apply brakes in this case on context sensitive elision and I'm a -1 here.

[mehdi_amini]

In D107236#2932063, @bondhugula wrote:

In D107236#2924729, @mehdi_amini wrote:

Does this really improve readability? Doesn't it create ambiguity for a reader as to which dialect's operation is being referring to (when you have the same names)? I might have misunderstood it but that's what the commit summary conveys to me.

Yes you’re right. I think this is about how it gets used though.
I’d claim that it is less intrusive than the custom assembly that we allow (we actually had internal teams complaining about it: affine, linalg, scf, … all have different custom conventions).

If you remember the original affine functions, I don’t think operations were prefixed: it was later during generalization that we added the prefix. The prefix wasn’t for readability but for necessary disambiguating in the parser.

The idea here is that you can see it like an extension of custom assembly: instead of just controlling the immediate operation you can impact the immediately nested regions in a limited way.

My main motivation is the TensorFlow graph dialect right now: when I write tfg.graph { I’d like to avoid prefixing every single operation with tfg. inside the graph itself.
This isn’t a region that allows anything else anyway, but even if it did, the few outlier operation would be prefixed.

Then, the ambiguity that this adds due to the prefix being dropped based on the context appears really undesirable to me at first sight. It's more readable with the prefix for me.

I think you'll have to elaborate with examples, because I don't see it really.

This also adds a layer of non-uniformity and surprises for people working across dialects and when folks encounter a new dialect. MLIR being an IR and not a user-facing programming language, I'm not sure you gain much by dropping a prefix like tfg to your ops.

The exact same argument applies to the custom assembly used by most dialect I believe.

Could you discuss this change on discourse? (Sorry if you already did and I missed it.)

Will do!

[mehdi_amini]

Will do!

https://llvm.discourse.group/t/rfc-allowing-to-set-a-default-dialect-to-be-used-for-printing-parsing-regions/4039

[GMNGeoffrey]

mlir/test/lib/Dialect/Test/TestOps.td

637–638

Comment is copy-pasta :-)

[lattner]

mlir/lib/IR/AsmPrinter.cpp

2310

We discussed a little offline, but I think it would be better to just remove the constraint that custom op printers print the operation name at the beginning. I don't see a real reason for them to do that, and it would also match the behavior of the parser (which doesn't parse the operation name, for likely obvious reasons).

I agree completely, the stream approach here is too-clever. It is also a pain (and pointless) to print the op name anyway, it would be far better to have the asmprinter do this by default.

[mehdi_amini]

Rebase on top of D108804 to remove the "filtering stream" now that the framework has direct control on printing the op name

[mehdi_amini]

Address Geoff's comment

[rriddle]

mlir/include/mlir/IR/OpAsmInterface.td

56

58

59–61

I don't think last statement is up-to-date. An empty string means no elision right? not that it is inherited.

mlir/lib/IR/AsmPrinter.cpp

2428

2671–2678

You could also use llvm::make_scope_exit instead.

2682

Shouldn't we push an empty string here instead?

mlir/lib/IR/Operation.cpp

651

Can we add a TODO/FIXME to kill this line?

mlir/lib/Parser/Parser.cpp

1892

Same comment here: why do we use builtin as the fallback here? I would expect an empty default dialect to mean "no ellision".

1903–1910

You can use llvm::make_scope_exit here instead.

[mehdi_amini]

Address River's comment

[mehdi_amini]

mlir/lib/IR/AsmPrinter.cpp

2671–2678

Grahh: believe me I searched for this utility and couldn't find it and concluded that my memory of its existence was bogus...

[rriddle]

Looks great, thanks!

mlir/include/mlir/IR/BuiltinOps.td

163–167

216–220

mlir/lib/IR/AsmPrinter.cpp

375

Parameter comments should have a =

2522–2523

?

mlir/test/lib/Dialect/Test/TestOps.td

637

[mehdi_amini]

Address River's comment

[silvas]

mlir/include/mlir/IR/OpImplementation.h

973 ↗

(On Diff #363584)

FYI, for future reference, this change was disruptive for us -- we had already updated tests to use builtin.func (in CHECK lines, etc.) or written new tests that were natively checking for builtin.func. After this change, we had to change it back.

[silvas]

mlir/include/mlir/IR/OpImplementation.h

973 ↗

(On Diff #363584)

(not sure if it was exactly this change, but something caused builtin.func to print as just func since our last integrate.)

[mehdi_amini]

mlir/include/mlir/IR/OpImplementation.h

973 ↗

(On Diff #363584)

This patch is the culprit indeed.
Something I'm not sure, is why was so many test changed to CHECK on builtin.func a few weeks ago when they could have continued to check for func alone? (I had to update many tests inside Google as well).

Before: you could parse without the builtin prefix but it would get printed, while now it'll be omitted if it can be parsed basically.

[silvas]

mlir/include/mlir/IR/OpImplementation.h

973 ↗

(On Diff #363584)

I think most tests are written by generate-test-checks.py which isn't smart enough to know that ;)

mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td

 include "mlir/IR/OpAsmInterface.td"
 include "mlir/IR/SymbolInterfaces.td"
 include "mlir/Interfaces/CallInterfaces.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
 // -----
 
 def SPV_AddressOfOp : SPV_Op<"mlir.addressof",
-    [DeclareOpInterfaceMethods<OpAsmOpInterface>, InFunctionScope, NoSideEffect]> {
+    [DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>,
+     InFunctionScope, NoSideEffect]> {
   let summary = "Get the address of a global variable.";
 
   let description = [{
     Variables in module scope are defined using symbol names. This op generates
     an SSA value that can be used to refer to the symbol within function scope
     for use in ops that expect an SSA value. This operation has no corresponding
     SPIR-V instruction; it's merely used for modelling purpose in the SPIR-V
     dialect. Since variables in module scope in SPIR-V dialect are of pointer
def SPV_AddressOfOp : SPV_Op<"mlir.addressof",
   let builders = [OpBuilder<(ins "spirv::GlobalVariableOp":$var)>];
 
   let assemblyFormat = "$variable attr-dict `:` type($pointer)";
 }
 
 // -----
 
 def SPV_ConstantOp : SPV_Op<"Constant",
-    [ConstantLike, DeclareOpInterfaceMethods<OpAsmOpInterface>, NoSideEffect]> {
+    [ConstantLike,
+     DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>,
+     NoSideEffect]> {
   let summary = "The op that declares a SPIR-V normal constant";
 
   let description = [{
     This op declares a SPIR-V normal constant. SPIR-V has multiple constant
     instructions covering different constant types:
 
     * `OpConstantTrue` and `OpConstantFalse` for boolean constants
     * `OpConstant` for scalar constants

mlir/include/mlir/Dialect/Shape/IR/ShapeOps.td

let extraClassDeclaration = [{
     // InferTypeOpInterface:
     static bool isCompatibleReturnTypes(TypeRange l, TypeRange r);
   }];
 }
 
 def Shape_ConstSizeOp : Shape_Op<"const_size", [
     ConstantLike,
     NoSideEffect,
-    DeclareOpInterfaceMethods<OpAsmOpInterface>
+    DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>
   ]> {
   let summary = "Creates a constant of type `shape.size`";
   let description = [{
     Creates a `shape.size` type representing the constant size given by `value`.
 
     ```mlir
     %x = shape.const_size 10
     ```

mlir/include/mlir/Dialect/StandardOps/IR/Ops.td

def CondBranchOp : Std_Op<"cond_br",
   }];
 }
 
 //===----------------------------------------------------------------------===//
 // ConstantOp
 //===----------------------------------------------------------------------===//
 
 def ConstantOp : Std_Op<"constant",
-    [ConstantLike, NoSideEffect, DeclareOpInterfaceMethods<OpAsmOpInterface>]> {
+    [ConstantLike, NoSideEffect,
+     DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>]> {
   let summary = "constant";
   let description = [{
     Syntax:
 
     ```
     operation ::= ssa-id `=` `std.constant` attribute-value `:` type
     ```
 

mlir/include/mlir/IR/BuiltinOps.h

 //===- BuiltinOps.h - MLIR Builtin Operations -------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the Builtin dialect's operations.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_IR_BUILTINOPS_H_
 #define MLIR_IR_BUILTINOPS_H_
 
 #include "mlir/IR/FunctionSupport.h"
+#include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/OwningOpRef.h"
 #include "mlir/IR/RegionKindInterface.h"
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Interfaces/CallInterfaces.h"
 #include "mlir/Interfaces/CastInterfaces.h"
 #include "mlir/Interfaces/DataLayoutInterfaces.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 #include "llvm/Support/PointerLikeTypeTraits.h"

mlir/include/mlir/IR/BuiltinOps.td

 // necessary for the validity of and defining the IR.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef BUILTIN_OPS
 #define BUILTIN_OPS
 
 include "mlir/IR/BuiltinDialect.td"
+include "mlir/IR/OpAsmInterface.td"
 include "mlir/IR/RegionKindInterface.td"
 include "mlir/IR/SymbolInterfaces.td"
 include "mlir/Interfaces/CallInterfaces.td"
 include "mlir/Interfaces/CastInterfaces.td"
 include "mlir/Interfaces/DataLayoutInterfaces.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
 // Base class for Builtin dialect ops.
]> {
   let printer = [{ return ::print(*this, p); }];
   let verifier = [{ return ::verify(*this); }];
 }
 
 //===----------------------------------------------------------------------===//
 // ModuleOp
 //===----------------------------------------------------------------------===//
 
 def ModuleOp : Builtin_Op<"module", [
-  AffineScope, IsolatedFromAbove, NoRegionArguments, SymbolTable, Symbol] 
-  # GraphRegionNoTerminator.traits> {
+    AffineScope, IsolatedFromAbove, NoRegionArguments, SymbolTable, Symbol,
+    OpAsmOpInterface
+  ] # GraphRegionNoTerminator.traits> {
   let summary = "A top level container operation";
   let description = [{
     A `module` represents a top-level container operation. It contains a single
     [graph region](../LangRef.md#control-flow-and-ssacfg-regions) containing a single block
     which can contain any operations and does not have a terminator. Operations
     within this region cannot implicitly capture values defined outside the module,
     i.e. Modules are [IsolatedFromAbove](../Traits.md#isolatedfromabove). Modules have
     an optional [symbol name](../SymbolsAndSymbolTables.md) which can be used to refer
     to them in operations.
let extraClassDeclaration = [{
     /// A ModuleOp may optionally define a symbol.
     bool isOptionalSymbol() { return true; }
 
     //===------------------------------------------------------------------===//
     // DataLayoutOpInterface Methods
     //===------------------------------------------------------------------===//
 
     DataLayoutSpecInterface getDataLayoutSpec();
+
+    //===------------------------------------------------------------------===//
+    // OpAsmOpInterface Methods
+    //===------------------------------------------------------------------===//
+
+    static ::llvm::StringRef getDefaultDialect() {
+      return "builtin";
+    }
   }];
   let verifier = [{ return ::verify(*this); }];
 
   // We need to ensure the block inside the region is properly terminated;
   // the auto-generated builders do not guarantee that.
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // UnrealizedConversionCastOp

mlir/include/mlir/IR/OpAsmInterface.td

InterfaceMethod<[{
         ```
 
         This would print the operation as follows:
 
         ```mlir
           %first_result, %middle_results:2, %0 = "my.op" ...
         ```
       }],
-      "void", "getAsmResultNames", (ins "::mlir::OpAsmSetValueNameFn":$setNameFn)
+      "void", "getAsmResultNames",
+      (ins "::mlir::OpAsmSetValueNameFn":$setNameFn),
+      "", ";"
+    >,
+    StaticInterfaceMethod<[{
+      Return the default dialect used when printing/parsing operations in
+      regions nested under this operation. This allows for eliding the dialect
+      prefix from the operation name, for example it would be possible to omit
+      the `spv.` prefix from all operations within a SpirV module if this method
+      returned `spv`. The default implementation returns an empty string which
+      is ignored.
+      }],
+      "StringRef", "getDefaultDialect", (ins), "", "return \"\";"
     >,
   ];
 }
 
 #endif // MLIR_OPASMINTERFACE

mlir/include/mlir/IR/OpDefinition.h

protected:
 
   /// Unless overridden, the custom assembly form of an op is always rejected.
   /// Op implementations should implement this to return failure.
   /// On success, they should fill in result with the fields to use.
   static ParseResult parse(OpAsmParser &parser, OperationState &result);
 
   // The fallback for the printer is to print it the generic assembly form.
   static void print(Operation *op, OpAsmPrinter &p);
-  static void printOpName(Operation *op, OpAsmPrinter &p);
+  static void printOpName(Operation *op, OpAsmPrinter &p,
+                          StringRef defaultDialect);
 
   /// Mutability management is handled by the OpWrapper/OpConstWrapper classes,
   /// so we can cast it away here.
   explicit OpState(Operation *state) : state(state) {}
 
 private:
   Operation *state;
 
static AbstractOperation::PrintAssemblyFn getPrintAssemblyFn() {
     return getPrintAssemblyFnImpl<ConcreteType>();
   }
   /// The internal implementation of `getPrintAssemblyFn` that is invoked when
   /// the concrete operation does not define a `print` method.
   template <typename ConcreteOpT>
   static std::enable_if_t<!detect_has_print<ConcreteOpT>::value,
                           AbstractOperation::PrintAssemblyFn>
   getPrintAssemblyFnImpl() {
-    return [](Operation *op, OpAsmPrinter &printer) {
+    return [](Operation *op, OpAsmPrinter &printer, StringRef defaultDialect) {
       return OpState::print(op, printer);
     };
   }
   /// The internal implementation of `getPrintAssemblyFn` that is invoked when
   /// the concrete operation defines a `print` method.
   template <typename ConcreteOpT>
   static std::enable_if_t<detect_has_print<ConcreteOpT>::value,
                           AbstractOperation::PrintAssemblyFn>
   getPrintAssemblyFnImpl() {
     return &printAssembly;
   }
-  static void printAssembly(Operation *op, OpAsmPrinter &p) {
-    OpState::printOpName(op, p);
+  static void printAssembly(Operation *op, OpAsmPrinter &p,
+                            StringRef defaultDialect) {
+    OpState::printOpName(op, p, defaultDialect);
     return cast<ConcreteType>(op).print(p);
   }
   /// Implementation of `VerifyInvariantsFn` AbstractOperation hook.
   static AbstractOperation::VerifyInvariantsFn getVerifyInvariantsFn() {
     return &verifyInvariants;
   }
 
   static constexpr bool hasNoDataMembers() {

mlir/include/mlir/IR/OperationSupport.h

public:
   using GetCanonicalizationPatternsFn =
       llvm::unique_function<void(RewritePatternSet &, MLIRContext *) const>;
   using FoldHookFn = llvm::unique_function<LogicalResult(
       Operation *, ArrayRef<Attribute>, SmallVectorImpl<OpFoldResult> &) const>;
   using HasTraitFn = llvm::unique_function<bool(TypeID) const>;
   using ParseAssemblyFn =
       llvm::unique_function<ParseResult(OpAsmParser &, OperationState &) const>;
   using PrintAssemblyFn =
-      llvm::unique_function<void(Operation *, OpAsmPrinter &) const>;
+      llvm::unique_function<void(Operation *, OpAsmPrinter &, StringRef) const>;
   using VerifyInvariantsFn =
       llvm::unique_function<LogicalResult(Operation *) const>;
 
   /// This is the name of the operation.
   const Identifier name;
 
   /// This is the dialect that this operation belongs to.
   Dialect &dialect;
 
   /// The unique identifier of the derived Op class.
   TypeID typeID;
 
   /// Use the specified object to parse this ops custom assembly format.
   ParseResult parseAssembly(OpAsmParser &parser, OperationState &result) const;
 
   /// Return the static hook for parsing this operation assembly.
   const ParseAssemblyFn &getParseAssemblyFn() const { return parseAssemblyFn; }
 
   /// This hook implements the AsmPrinter for this operation.
-  void printAssembly(Operation *op, OpAsmPrinter &p) const {
-    return printAssemblyFn(op, p);
+  void printAssembly(Operation *op, OpAsmPrinter &p,
+                     StringRef defaultDialect) const {
+    return printAssemblyFn(op, p, defaultDialect);
   }
 
   /// This hook implements the verifier for this operation.  It should emits an
   /// error message and returns failure if a problem is detected, or returns
   /// success if everything is ok.
   LogicalResult verifyInvariants(Operation *op) const {
     return verifyInvariantsFn(op);
   }

mlir/lib/IR/AsmPrinter.cpp

 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/Operation.h"
 #include "mlir/IR/SubElementInterfaces.h"
 #include "llvm/ADT/APFloat.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/MapVector.h"
 #include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/ScopeExit.h"
 #include "llvm/ADT/ScopedHashTable.h"
 #include "llvm/ADT/SetVector.h"
 #include "llvm/ADT/SmallString.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/StringSet.h"
 #include "llvm/ADT/TypeSwitch.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Endian.h"
void print(Operation *op) {
     if (printerFlags.shouldPrintDebugInfo())
       initializer.visit(op->getLoc(), /*canBeDeferred=*/true);
 
     // If requested, always print the generic form.
     if (!printerFlags.shouldPrintGenericOpForm()) {
       // Check to see if this is a known operation.  If so, use the registered
       // custom printer hook.
       if (auto *opInfo = op->getAbstractOperation()) {
-        opInfo->printAssembly(op, *this);
+        opInfo->printAssembly(op, *this, /*defaultDialect=*/"");
         return;
       }
     }
 
     // Otherwise print with the generic assembly form.
     printGenericOp(op);
   }
 

 //===----------------------------------------------------------------------===//
 // OperationPrinter
 //===----------------------------------------------------------------------===//
 
 namespace {
 /// This class contains the logic for printing operations, regions, and blocks.
 class OperationPrinter : public ModulePrinter, private OpAsmPrinter {
 public:
   explicit OperationPrinter(raw_ostream &os, OpPrintingFlags flags,
                             AsmStateImpl &state)
       : ModulePrinter(os, flags, &state) {}
 
   /// Print the given top-level operation.
   void printTopLevelOperation(Operation *op);
 
   /// Print the given operation with its indent and location.
   void print(Operation *op);
   /// Print the bare location, not including indentation/location/etc.
   void printOperation(Operation *op);
   /// Print the given operation in the generic form.
   void printGenericOp(Operation *op, bool printOpName) override;
 
   /// Print the name of the given block.
   void printBlockName(Block *block);
 
   /// Print the given block. If 'printBlockArgs' is false, the arguments of the
   /// block are not printed. If 'printBlockTerminator' is false, the terminator
   /// operation of the block is not printed.
   void print(Block *block, bool printBlockArgs = true,
              bool printBlockTerminator = true);
 
   /// Print the ID of the given value, optionally with its result number.
   void printValueID(Value value, bool printResultNo = true,
                     raw_ostream *streamOverride = nullptr) const;
 
   //===--------------------------------------------------------------------===//
   // OpAsmPrinter methods
   //===--------------------------------------------------------------------===//
 
   /// Return the current stream of the printer.
   raw_ostream &getStream() const override { return os; }
 
   /// Print a newline and indent the printer to the start of the current
   /// operation.
   void printNewline() override {
     os << newLine;
     os.indent(currentIndent);
   }
 
   /// Print the given type.
   void printType(Type type) override { ModulePrinter::printType(type); }
 
   /// Print the given attribute.
   void printAttribute(Attribute attr) override {
     ModulePrinter::printAttribute(attr);
   }
 
   /// Print the given attribute without its type. The corresponding parser must
   /// provide a valid type for the attribute.
   void printAttributeWithoutType(Attribute attr) override {
     ModulePrinter::printAttribute(attr, AttrTypeElision::Must);
   }
 
   /// Print a block argument in the usual format of:
   ///   %ssaName : type {attr1=42} loc("here")
   /// where location printing is controlled by the standard internal option.
   /// You may pass omitType=true to not print a type, and pass an empty
   /// attribute list if you don't care for attributes.
void printAffineExprOfSSAIds(AffineExpr expr, ValueRange dimOperands,
                                ValueRange symOperands) override;
 
   /// Print the given string as a symbol reference.
   void printSymbolName(StringRef symbolRef) override {
     ::printSymbolReference(symbolRef, os);
   }
 
 private:
+  // Contains the stack of default dialects to use when printing regions.
+  // A new dialect is pushed to the stack before parsing regions nested under an
+  // operation implementing `OpAsmOpInterface`, and popped when done. At the
+  // top-level we start with "builtin" as the default, so that the top-level
+  // `module` operation prints as-is.
+  SmallVector<StringRef> defaultDialectStack{"builtin"};
+
   /// The number of spaces used for indenting nested operations.
   const static unsigned indentWidth = 2;
 
   // This is the current indentation level for nested structures.
   unsigned currentIndent = 0;
 };
 } // end anonymous namespace
 
if (size_t numResults = op->getNumResults()) {
     os << " = ";
   }
 
   // If requested, always print the generic form.
   if (!printerFlags.shouldPrintGenericOpForm()) {
     // Check to see if this is a known operation.  If so, use the registered
     // custom printer hook.
     if (auto *opInfo = op->getAbstractOperation()) {
-      opInfo->printAssembly(op, *this);
+      opInfo->printAssembly(op, *this, defaultDialectStack.back());
       return;
     }
     // Otherwise try to dispatch to the dialect, if available.
     if (Dialect *dialect = op->getDialect()) {
       if (auto opPrinter = dialect->getOperationPrinter(op)) {
         // Print the op name first.
         StringRef name = op->getName().getStringRef();
+        name.consume_front((defaultDialectStack.back() + ".").str());
         printEscapedString(name, os);
         // Print the rest of the op now.
         opPrinter(op, *this);
         return;
       }
     }
   }
 
   // Otherwise print with the generic assembly form.
void OperationPrinter::printSuccessorAndUseList(Block *successor,
   os << ')';
 }
 
 void OperationPrinter::printRegion(Region &region, bool printEntryBlockArgs,
                                    bool printBlockTerminators,
                                    bool printEmptyBlock) {
   os << " {" << newLine;
   if (!region.empty()) {
+    auto restoreDefaultDialect =
+        llvm::make_scope_exit([&]() { defaultDialectStack.pop_back(); });
+    if (auto iface = dyn_cast<OpAsmOpInterface>(region.getParentOp()))
+      defaultDialectStack.push_back(iface.getDefaultDialect());
+    else
+      defaultDialectStack.push_back("");
+
     auto *entryBlock = &region.front();
     // Force printing the block header if printEmptyBlock is set and the block
     // is empty or if printEntryBlockArgs is set and there are arguments to
     // print.
     bool shouldAlwaysPrintBlockHeader =
         (printEmptyBlock && entryBlock->empty()) ||
         (printEntryBlockArgs && entryBlock->getNumArguments() != 0);
     print(entryBlock, shouldAlwaysPrintBlockHeader, printBlockTerminators);
     for (auto &b : llvm::drop_begin(region.getBlocks(), 1))
       print(&b);
   }
   os.indent(currentIndent) << "}";
 }
 
 void OperationPrinter::printAffineMapOfSSAIds(AffineMapAttr mapAttr,

mlir/lib/IR/Operation.cpp

 // The fallback for the parser is to reject the custom assembly form.
 ParseResult OpState::parse(OpAsmParser &parser, OperationState &result) {
   return parser.emitError(parser.getNameLoc(), "has no custom assembly form");
 }
 
 // The fallback for the printer is to print in the generic assembly form.
 void OpState::print(Operation *op, OpAsmPrinter &p) { p.printGenericOp(op); }
 // The fallback for the printer is to print in the generic assembly form.
-void OpState::printOpName(Operation *op, OpAsmPrinter &p) {
+void OpState::printOpName(Operation *op, OpAsmPrinter &p,
+                          StringRef defaultDialect) {
   StringRef name = op->getName().getStringRef();
-  if (name.startswith("std."))
+  if (name.startswith((defaultDialect + ".").str()))
+    name = name.drop_front(defaultDialect.size() + 1);
+  // TODO: remove this special case.
+  else if (name.startswith("std."))
     name = name.drop_front(4);
   p.getStream() << name;
 }
 
 /// Emit an error about fatal conditions with this operation, reporting up to
 /// any diagnostic handlers that may be listening.
 InFlightDiagnostic OpState::emitError(const Twine &message) {
   return getOperation()->emitError(message);

mlir/lib/Parser/Parser.cpp

 #include "Parser.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/Verifier.h"
 #include "mlir/Parser.h"
 #include "mlir/Parser/AsmParserState.h"
 #include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/ScopeExit.h"
 #include "llvm/ADT/StringSet.h"
 #include "llvm/ADT/bit.h"
 #include "llvm/Support/PrettyStackTrace.h"
 #include "llvm/Support/SourceMgr.h"
 #include <algorithm>
 
 using namespace mlir;
 using namespace mlir::detail;
private:
   /// A flag that indicates if any errors were emitted during parsing.
   bool emittedError = false;
 };
 } // end anonymous namespace.
 
 Operation *
 OperationParser::parseCustomOperation(ArrayRef<ResultRecord> resultIDs) {
   llvm::SMLoc opLoc = getToken().getLoc();
-  StringRef opName = getTokenSpelling();
+  std::string opName = getTokenSpelling().str();
   auto *opDefinition = AbstractOperation::lookup(opName, getContext());
+  StringRef defaultDialect = getState().defaultDialectStack.back();
   Dialect *dialect = nullptr;
   if (opDefinition) {
     dialect = &opDefinition->dialect;
   } else {
-    if (opName.contains('.')) {
+    if (StringRef(opName).contains('.')) {
       // This op has a dialect, we try to check if we can register it in the
       // context on the fly.
-      StringRef dialectName = opName.split('.').first;
+      StringRef dialectName = StringRef(opName).split('.').first;
       dialect = getContext()->getLoadedDialect(dialectName);
       if (!dialect && (dialect = getContext()->getOrLoadDialect(dialectName)))
         opDefinition = AbstractOperation::lookup(opName, getContext());
     } else {
-      // If the operation name has no namespace prefix we treat it as a builtin
-      // or standard operation and prefix it with "builtin" or "std".
-      // TODO: Remove the special casing here.
-      opDefinition = AbstractOperation::lookup(Twine("builtin." + opName).str(),
-                                               getContext());
+      // If the operation name has no namespace prefix we lookup the current
+      // default dialect (set through OpAsmOpInterface).
+      opDefinition = AbstractOperation::lookup(
+          Twine(defaultDialect + "." + opName).str(), getContext());
       if (!opDefinition && getContext()->getOrLoadDialect("std")) {
         opDefinition = AbstractOperation::lookup(Twine("std." + opName).str(),
                                                  getContext());
       }
-      if (opDefinition)
-        opName = opDefinition->name.strref();
+      if (opDefinition) {
+        dialect = &opDefinition->dialect;
+        opName = opDefinition->name.str();
+      } else if (!defaultDialect.empty()) {
+        dialect = getContext()->getOrLoadDialect(defaultDialect);
+        opName = (defaultDialect + "." + opName).str();
+      }
     }
   }
 
   // This is the actual hook for the custom op parsing, usually implemented by
   // the op itself (`Op::parse()`). We retrieve it either from the
   // AbstractOperation or from the Dialect.
   function_ref<ParseResult(OpAsmParser &, OperationState &)> parseAssemblyFn;
   bool isIsolatedFromAbove = false;
 
+  defaultDialect = "";
   if (opDefinition) {
     parseAssemblyFn = opDefinition->getParseAssemblyFn();
     isIsolatedFromAbove =
         opDefinition->hasTrait<OpTrait::IsIsolatedFromAbove>();
+    auto *iface = opDefinition->getInterface<OpAsmOpInterface>();
+    if (iface && !iface->getDefaultDialect().empty())
+      defaultDialect = iface->getDefaultDialect();
   } else {
     Optional<Dialect::ParseOpHook> dialectHook;
     if (dialect)
       dialectHook = dialect->getParseOperationHook(opName);
     if (!dialectHook.hasValue()) {
       emitError(opLoc) << "custom op '" << opName << "' is unknown";
       return nullptr;
     }
     parseAssemblyFn = *dialectHook;
   }
+  getState().defaultDialectStack.push_back(defaultDialect);
+  auto restoreDefaultDialect = llvm::make_scope_exit(
+      [&]() { getState().defaultDialectStack.pop_back(); });
 
   consumeToken();
 
   // If the custom op parser crashes, produce some indication to help
   // debugging.
-  std::string opNameStr = opName.str();
   llvm::PrettyStackTraceFormat fmt("MLIR Parser: custom op parser '%s'",
-                                   opNameStr.c_str());
+                                   opName.c_str());
 
   // Get location information for the operation.
   auto srcLocation = getEncodedSourceLocation(opLoc);
   OperationState opState(srcLocation, opName);
 
   // If we are populating the parser state, start a new operation definition.
   if (state.asmState)
     state.asmState->startOperationDefinition(opState.name);

mlir/lib/Parser/ParserState.h

struct ParserState {
   SymbolState &symbols;
 
   /// The depth of this parser in the nested parsing stack.
   size_t parserDepth;
 
   /// An optional pointer to a struct containing high level parser state to be
   /// populated during parsing.
   AsmParserState *asmState;
+
+  // Contains the stack of default dialect to use when parsing regions.
+  // A new dialect get pushed to the stack before parsing regions nested
+  // under an operation implementing `OpAsmOpInterface`, and
+  // popped when done. At the top-level we start with "builtin" as the
+  // default, so that the top-level `module` operation parses as-is.
+  SmallVector<StringRef> defaultDialectStack{"builtin"};
 };
 
 } // end namespace detail
 } // end namespace mlir
 
 #endif // MLIR_LIB_PARSER_PARSERSTATE_H

mlir/test/Analysis/test-shape-fn-report.mlir

func @tanh(%arg: tensor<10x20xf32>) -> tensor<10x20xf32>
   // expected-remark@+1 {{associated shape function: same_result_shape}}
   %1 = "test.same_operand_result_type"(%0) : (tensor<10x20xf32>) -> tensor<10x20xf32>
   return %1 : tensor<10x20xf32>
 }
 
 // The shape function library with some local functions.
 shape.function_library @shape_lib {
   // Test shape function that returns the shape of input arg as result shape.
-  func @same_result_shape(%arg: !shape.value_shape) -> !shape.shape {
+  builtin.func @same_result_shape(%arg: !shape.value_shape) -> !shape.shape {
     %0 = shape.shape_of %arg : !shape.value_shape -> !shape.shape
     return %0 : !shape.shape
   }
 } mapping {
   test.same_operand_result_type = @same_result_shape
 }
 
 }

mlir/test/Conversion/GPUToNVVM/gpu-to-nvvm.mlir

 // RUN: mlir-opt %s -convert-gpu-to-nvvm -split-input-file | FileCheck %s
 // RUN: mlir-opt %s -convert-gpu-to-nvvm='index-bitwidth=32' -split-input-file | FileCheck --check-prefix=CHECK32 %s
 
 gpu.module @test_module {
   // CHECK-LABEL: func @gpu_index_ops()
   // CHECK32-LABEL: func @gpu_index_ops()
-  func @gpu_index_ops()
+  builtin.func @gpu_index_ops()
       -> (index, index, index, index, index, index,
           index, index, index, index, index, index) {
     // CHECK32-NOT: = llvm.sext %{{.*}} : i32 to i64
 
     // CHECK: = nvvm.read.ptx.sreg.tid.x : i32
     // CHECK: = llvm.sext %{{.*}} : i32 to i64
     %tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index)
     // CHECK: = nvvm.read.ptx.sreg.tid.y : i32
gpu.module @test_module {
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK-LABEL: func @gpu_index_comp
   // CHECK32-LABEL: func @gpu_index_comp
-  func @gpu_index_comp(%idx : index) -> index {
+  builtin.func @gpu_index_comp(%idx : index) -> index {
     // CHECK: = llvm.add %{{.*}}, %{{.*}} : i64
     // CHECK32: = llvm.add %{{.*}}, %{{.*}} : i32
     %0 = addi %idx, %idx : index
     // CHECK: llvm.return %{{.*}} : i64
     // CHECK32: llvm.return %{{.*}} : i32
     std.return %0 : index
   }
 }
gpu.func @gpu_all_reduce_region() {
     gpu.return
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK-LABEL: func @gpu_shuffle()
-  func @gpu_shuffle() -> (f32) {
+  builtin.func @gpu_shuffle() -> (f32) {
     // CHECK: %[[#VALUE:]] = llvm.mlir.constant(1.000000e+00 : f32) : f32
     %arg0 = constant 1.0 : f32
     // CHECK: %[[#OFFSET:]] = llvm.mlir.constant(4 : i32) : i32
     %arg1 = constant 4 : i32
     // CHECK: %[[#WIDTH:]] = llvm.mlir.constant(23 : i32) : i32
     %arg2 = constant 23 : i32
     // CHECK: %[[#ONE:]] = llvm.mlir.constant(1 : i32) : i32
     // CHECK: %[[#SHL:]] = llvm.shl %[[#ONE]], %[[#WIDTH]] : i32
     // CHECK: %[[#MASK:]] = llvm.sub %[[#SHL]], %[[#ONE]] : i32
     // CHECK: %[[#CLAMP:]] = llvm.sub %[[#WIDTH]], %[[#ONE]] : i32
     // CHECK: %[[#SHFL:]] = nvvm.shfl.sync.bfly %[[#MASK]], %[[#VALUE]], %[[#OFFSET]], %[[#CLAMP]] : !llvm.struct<(f32, i1)>
     // CHECK: llvm.extractvalue %[[#SHFL]][0 : index] : !llvm.struct<(f32, i1)>
     // CHECK: llvm.extractvalue %[[#SHFL]][1 : index] : !llvm.struct<(f32, i1)>
     %shfl, %pred = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = "xor" } : (f32, i32, i32) -> (f32, i1)
 
     std.return %shfl : f32
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK-LABEL: func @gpu_sync()
-  func @gpu_sync() {
+  builtin.func @gpu_sync() {
     // CHECK: nvvm.barrier0
     gpu.barrier
     std.return
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_fabsf(f32) -> f32
   // CHECK: llvm.func @__nv_fabs(f64) -> f64
   // CHECK-LABEL: func @gpu_fabs
-  func @gpu_fabs(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_fabs(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = std.absf %arg_f32 : f32
     // CHECK: llvm.call @__nv_fabsf(%{{.*}}) : (f32) -> f32
     %result64 = std.absf %arg_f64 : f64
     // CHECK: llvm.call @__nv_fabs(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_ceilf(f32) -> f32
   // CHECK: llvm.func @__nv_ceil(f64) -> f64
   // CHECK-LABEL: func @gpu_ceil
-  func @gpu_ceil(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_ceil(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = std.ceilf %arg_f32 : f32
     // CHECK: llvm.call @__nv_ceilf(%{{.*}}) : (f32) -> f32
     %result64 = std.ceilf %arg_f64 : f64
     // CHECK: llvm.call @__nv_ceil(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_floorf(f32) -> f32
   // CHECK: llvm.func @__nv_floor(f64) -> f64
   // CHECK-LABEL: func @gpu_floor
-  func @gpu_floor(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_floor(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = std.floorf %arg_f32 : f32
     // CHECK: llvm.call @__nv_floorf(%{{.*}}) : (f32) -> f32
     %result64 = std.floorf %arg_f64 : f64
     // CHECK: llvm.call @__nv_floor(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_cosf(f32) -> f32
   // CHECK: llvm.func @__nv_cos(f64) -> f64
   // CHECK-LABEL: func @gpu_cos
-  func @gpu_cos(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_cos(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.cos %arg_f32 : f32
     // CHECK: llvm.call @__nv_cosf(%{{.*}}) : (f32) -> f32
     %result64 = math.cos %arg_f64 : f64
     // CHECK: llvm.call @__nv_cos(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_expf(f32) -> f32
   // CHECK: llvm.func @__nv_exp(f64) -> f64
   // CHECK-LABEL: func @gpu_exp
-  func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.exp %arg_f32 : f32
     // CHECK: llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32
     %result64 = math.exp %arg_f64 : f64
     // CHECK: llvm.call @__nv_exp(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_exp2f(f32) -> f32
   // CHECK: llvm.func @__nv_exp2(f64) -> f64
   // CHECK-LABEL: func @gpu_exp2
-  func @gpu_exp2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_exp2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.exp2 %arg_f32 : f32
     // CHECK: llvm.call @__nv_exp2f(%{{.*}}) : (f32) -> f32
     %result64 = math.exp2 %arg_f64 : f64
     // CHECK: llvm.call @__nv_exp2(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_logf(f32) -> f32
   // CHECK: llvm.func @__nv_log(f64) -> f64
   // CHECK-LABEL: func @gpu_log
-  func @gpu_log(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_log(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.log %arg_f32 : f32
     // CHECK: llvm.call @__nv_logf(%{{.*}}) : (f32) -> f32
     %result64 = math.log %arg_f64 : f64
     // CHECK: llvm.call @__nv_log(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_log10f(f32) -> f32
   // CHECK: llvm.func @__nv_log10(f64) -> f64
   // CHECK-LABEL: func @gpu_log10
-  func @gpu_log10(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_log10(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.log10 %arg_f32 : f32
     // CHECK: llvm.call @__nv_log10f(%{{.*}}) : (f32) -> f32
     %result64 = math.log10 %arg_f64 : f64
     // CHECK: llvm.call @__nv_log10(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_log1pf(f32) -> f32
   // CHECK: llvm.func @__nv_log1p(f64) -> f64
   // CHECK-LABEL: func @gpu_log1p
-  func @gpu_log1p(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_log1p(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.log1p %arg_f32 : f32
     // CHECK: llvm.call @__nv_log1pf(%{{.*}}) : (f32) -> f32
     %result64 = math.log1p %arg_f64 : f64
     // CHECK: llvm.call @__nv_log1p(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_log2f(f32) -> f32
   // CHECK: llvm.func @__nv_log2(f64) -> f64
   // CHECK-LABEL: func @gpu_log2
-  func @gpu_log2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_log2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.log2 %arg_f32 : f32
     // CHECK: llvm.call @__nv_log2f(%{{.*}}) : (f32) -> f32
     %result64 = math.log2 %arg_f64 : f64
     // CHECK: llvm.call @__nv_log2(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_sinf(f32) -> f32
   // CHECK: llvm.func @__nv_sin(f64) -> f64
   // CHECK-LABEL: func @gpu_sin
-  func @gpu_sin(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_sin(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.sin %arg_f32 : f32
     // CHECK: llvm.call @__nv_sinf(%{{.*}}) : (f32) -> f32
     %result64 = math.sin %arg_f64 : f64
     // CHECK: llvm.call @__nv_sin(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_tanhf(f32) -> f32
   // CHECK: llvm.func @__nv_tanh(f64) -> f64
   // CHECK-LABEL: func @gpu_tanh
-  func @gpu_tanh(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) -> (f16, f32, f64) {
+  builtin.func @gpu_tanh(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) -> (f16, f32, f64) {
     %result16 = math.tanh %arg_f16 : f16
     // CHECK: llvm.fpext %{{.*}} : f16 to f32
     // CHECK-NEXT: llvm.call @__nv_tanhf(%{{.*}}) : (f32) -> f32
     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
     %result32 = math.tanh %arg_f32 : f32
     // CHECK: llvm.call @__nv_tanhf(%{{.*}}) : (f32) -> f32
     %result64 = math.tanh %arg_f64 : f64
     // CHECK: llvm.call @__nv_tanh(%{{.*}}) : (f64) -> f64
     std.return %result16, %result32, %result64 : f16, f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_rsqrtf(f32) -> f32
   // CHECK: llvm.func @__nv_rsqrt(f64) -> f64
   // CHECK-LABEL: func @gpu_rsqrt
-  func @gpu_rsqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
+  builtin.func @gpu_rsqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
       -> (f16, f32, f64) {
     %result16 = math.rsqrt %arg_f16 : f16
     // CHECK: llvm.fpext %{{.*}} : f16 to f32
     // CHECK-NEXT: llvm.call @__nv_rsqrtf(%{{.*}}) : (f32) -> f32
     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
     %result32 = math.rsqrt %arg_f32 : f32
     // CHECK: llvm.call @__nv_rsqrtf(%{{.*}}) : (f32) -> f32
     %result64 = math.rsqrt %arg_f64 : f64
     // CHECK: llvm.call @__nv_rsqrt(%{{.*}}) : (f64) -> f64
     std.return %result16, %result32, %result64 : f16, f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_sqrtf(f32) -> f32
   // CHECK: llvm.func @__nv_sqrt(f64) -> f64
   // CHECK-LABEL: func @gpu_sqrt
-  func @gpu_sqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
+  builtin.func @gpu_sqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
       -> (f16, f32, f64) {
     %result16 = math.sqrt %arg_f16 : f16
     // CHECK: llvm.fpext %{{.*}} : f16 to f32
     // CHECK-NEXT: llvm.call @__nv_sqrtf(%{{.*}}) : (f32) -> f32
     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
     %result32 = math.sqrt %arg_f32 : f32
     // CHECK: llvm.call @__nv_sqrtf(%{{.*}}) : (f32) -> f32
     %result64 = math.sqrt %arg_f64 : f64
     // CHECK: llvm.call @__nv_sqrt(%{{.*}}) : (f64) -> f64
     std.return %result16, %result32, %result64 : f16, f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_atanf(f32) -> f32
   // CHECK: llvm.func @__nv_atan(f64) -> f64
   // CHECK-LABEL: func @gpu_atan
-  func @gpu_atan(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
+  builtin.func @gpu_atan(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
       -> (f16, f32, f64) {
     %result16 = math.atan %arg_f16 : f16
     // CHECK: llvm.fpext %{{.*}} : f16 to f32
     // CHECK-NEXT: llvm.call @__nv_atanf(%{{.*}}) : (f32) -> f32
     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
     %result32 = math.atan %arg_f32 : f32
     // CHECK: llvm.call @__nv_atanf(%{{.*}}) : (f32) -> f32
     %result64 = math.atan %arg_f64 : f64
     // CHECK: llvm.call @__nv_atan(%{{.*}}) : (f64) -> f64
     std.return %result16, %result32, %result64 : f16, f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_atan2f(f32, f32) -> f32
   // CHECK: llvm.func @__nv_atan2(f64, f64) -> f64
   // CHECK-LABEL: func @gpu_atan2
-  func @gpu_atan2(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
+  builtin.func @gpu_atan2(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
       -> (f16, f32, f64) {
     %result16 = math.atan2 %arg_f16, %arg_f16 : f16
     // CHECK: llvm.fpext %{{.*}} : f16 to f32
     // CHECK: llvm.fpext %{{.*}} : f16 to f32
     // CHECK-NEXT: llvm.call @__nv_atan2f(%{{.*}}) : (f32, f32) -> f32
     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
     %result32 = math.atan2 %arg_f32, %arg_f32 : f32
     // CHECK: llvm.call @__nv_atan2f(%{{.*}}) : (f32, f32) -> f32
     %result64 = math.atan2 %arg_f64, %arg_f64 : f64
     // CHECK: llvm.call @__nv_atan2(%{{.*}}) : (f64, f64) -> f64
     std.return %result16, %result32, %result64 : f16, f32, f64
   }
 }
 
 // -----
 
 // Test that we handled properly operation with SymbolTable other than module op
 gpu.module @test_module {
   "test.symbol_scope"() ({
   // CHECK: test.symbol_scope
   // CHECK: llvm.func @__nv_expf(f32) -> f32
   // CHECK: llvm.func @__nv_exp(f64) -> f64
   // CHECK-LABEL: func @gpu_exp
-    func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+    builtin.func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
       %result32 = math.exp %arg_f32 : f32
       // CHECK: llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32
       %result64 = math.exp %arg_f64 : f64
       // CHECK: llvm.call @__nv_exp(%{{.*}}) : (f64) -> f64
       std.return %result32, %result64 : f32, f64
     }
     "test.finish" () : () -> ()
   }) : () -> ()
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_expm1f(f32) -> f32
   // CHECK: llvm.func @__nv_expm1(f64) -> f64
   // CHECK-LABEL: func @gpu_expm1
-  func @gpu_expm1(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_expm1(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.expm1 %arg_f32 : f32
     // CHECK: llvm.call @__nv_expm1f(%{{.*}}) : (f32) -> f32
     %result64 = math.expm1 %arg_f64 : f64
     // CHECK: llvm.call @__nv_expm1(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__nv_powf(f32, f32) -> f32
   // CHECK: llvm.func @__nv_pow(f64, f64) -> f64
   // CHECK-LABEL: func @gpu_pow
-  func @gpu_pow(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_pow(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.powf %arg_f32, %arg_f32 : f32
     // CHECK: llvm.call @__nv_powf(%{{.*}}, %{{.*}}) : (f32, f32) -> f32
     %result64 = math.powf %arg_f64, %arg_f64 : f64
     // CHECK: llvm.call @__nv_pow(%{{.*}}, %{{.*}}) : (f64, f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 

mlir/test/Conversion/GPUToNVVM/wmma-ops-to-nvvm.mlir

 // RUN: mlir-opt --convert-gpu-to-nvvm="index-bitwidth=32" --split-input-file %s | FileCheck %s
 
 gpu.module @test_module {
 
   // CHECK-LABEL: func @gpu_wmma_load_op() ->
   // CHECK-SAME: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> {
-  func @gpu_wmma_load_op() -> (!gpu.mma_matrix<16x16xf16, "AOp">) {
+  builtin.func @gpu_wmma_load_op() -> (!gpu.mma_matrix<16x16xf16, "AOp">) {
     %wg = memref.alloca() {alignment = 32} : memref<32x32xf16, 3>
     %i = constant 16 : index
     %j = constant 16 : index
     %0 = gpu.subgroup_mma_load_matrix %wg[%i, %j] {leadDimension = 32 : index} : memref<32x32xf16, 3> -> !gpu.mma_matrix<16x16xf16, "AOp">
     // CHECK:  %[[INX:.*]] = llvm.mlir.constant(16 : index) : i32
     // CHECK: %{{.*}} = llvm.insertvalue %{{.*}}, %{{.*}}[{{.*}}, {{.*}}]
     // CHECK:  %[[LDM:.*]] = llvm.mlir.constant(32 : index) : i32
     // CHECK:  %[[LI:.*]] = llvm.mul %[[LDM]], %[[INX]]  : i32

 }
 
 // -----
 
 gpu.module @test_module {
 
   // CHECK-LABEL: func @gpu_wmma_store_op
   // CHECK-SAME: (%[[D:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>) {
-  func @gpu_wmma_store_op(%arg0 : !gpu.mma_matrix<16x16xf16, "COp">) -> () {
+  builtin.func @gpu_wmma_store_op(%arg0 : !gpu.mma_matrix<16x16xf16, "COp">) -> () {
     %sg = memref.alloca(){alignment = 32} : memref<32x32xf16, 3>
     %i = constant 16 : index
     %j = constant 16 : index
     gpu.subgroup_mma_store_matrix %arg0, %sg[%i,%j] {leadDimension= 32 : index} : !gpu.mma_matrix<16x16xf16, "COp">, memref<32x32xf16, 3>
     // CHECK:  %[[INX:.*]] = llvm.mlir.constant(16 : index) : i32
     // CHECK: %{{.*}} = llvm.insertvalue %{{.*}}, %{{.*}}[{{.*}}, {{.*}}]
     // CHECK:  %[[LDM:.*]] = llvm.mlir.constant(32 : index) : i32
     // CHECK:  %[[LI:.*]] = llvm.mul %[[LDM]], %[[INX]]  : i32

 }
 
 // -----
 
 gpu.module @test_module {
 
   // CHECK-LABEL: func @gpu_wmma_mma_op
   // CHECK-SAME: (%[[A:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>, %[[B:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>, %[[C:.*]]: !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>)
-  func @gpu_wmma_mma_op(%A : !gpu.mma_matrix<16x16xf16, "AOp">, %B : !gpu.mma_matrix<16x16xf16, "BOp">, %C : !gpu.mma_matrix<16x16xf16, "COp">) -> (!gpu.mma_matrix<16x16xf16, "COp">) {
+  builtin.func @gpu_wmma_mma_op(%A : !gpu.mma_matrix<16x16xf16, "AOp">, %B : !gpu.mma_matrix<16x16xf16, "BOp">, %C : !gpu.mma_matrix<16x16xf16, "COp">) -> (!gpu.mma_matrix<16x16xf16, "COp">) {
     %D = gpu.subgroup_mma_compute %A, %B, %C : !gpu.mma_matrix<16x16xf16, "AOp">, !gpu.mma_matrix<16x16xf16, "BOp"> -> !gpu.mma_matrix<16x16xf16, "COp">
     // CHECK:  %[[A1:.*]] = llvm.extractvalue %[[A]][0 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
     // CHECK:  %[[A2:.*]] = llvm.extractvalue %[[A]][1 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
     // CHECK:  %[[A3:.*]] = llvm.extractvalue %[[A]][2 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
     // CHECK:  %[[A4:.*]] = llvm.extractvalue %[[A]][3 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
     // CHECK:  %[[A5:.*]] = llvm.extractvalue %[[A]][4 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
     // CHECK:  %[[A6:.*]] = llvm.extractvalue %[[A]][5 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
     // CHECK:  %[[A7:.*]] = llvm.extractvalue %[[A]][6 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>

 //       CHECK:   llvm.br ^bb1(%{{.*}}, %[[ACC_MUL]] : i32, !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>)
 //       CHECK:  ^bb3:  // pred: ^bb1
 //       CHECK:   %87 = llvm.extractvalue %[[ACC]][0 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK:   %88 = llvm.extractvalue %[[ACC]][1 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK:   %89 = llvm.extractvalue %[[ACC]][2 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK:   %90 = llvm.extractvalue %[[ACC]][3 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK:   nvvm.wmma.m16n16k16.store.d.f16.row.stride %86, %87, %88, %89, %90, %79 : !llvm.ptr<i32>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, i32
 
-  func @gpu_wmma_mma_loop_op(%arg0: memref<128x128xf16>, %arg1: memref<128x128xf16>, %arg2: memref<128x128xf16>) {
+  builtin.func @gpu_wmma_mma_loop_op(%arg0: memref<128x128xf16>, %arg1: memref<128x128xf16>, %arg2: memref<128x128xf16>) {
       %c0 = constant 0 : index
       %c128 = constant 128 : index
       %c32 = constant 32 : index
       %0 = gpu.subgroup_mma_load_matrix %arg2[%c0, %c0] {leadDimension = 128 : index} : memref<128x128xf16> -> !gpu.mma_matrix<16x16xf16, "COp">
       br ^bb1(%c0, %0 : index, !gpu.mma_matrix<16x16xf16, "COp">)
     ^bb1(%1: index, %2: !gpu.mma_matrix<16x16xf16, "COp">):  // 2 preds: ^bb0, ^bb2
       %3 = cmpi slt, %1, %c128 : index
       cond_br %3, ^bb2, ^bb3

 //       CHECK: %[[C1:.+]] = llvm.mlir.constant(1 : i32) : i32
 //       CHECK: %[[V2:.+]] = llvm.insertelement %[[CST]], %[[V1]][%[[C1]] : i32] : vector<2xf16>
 //       CHECK: %[[M0:.+]] = llvm.mlir.undef : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK: %[[M1:.+]] = llvm.insertvalue %[[V2]], %[[M0]][0 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK: %[[M2:.+]] = llvm.insertvalue %[[V2]], %[[M1]][1 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK: %[[M3:.+]] = llvm.insertvalue %[[V2]], %[[M2]][2 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK: %[[M4:.+]] = llvm.insertvalue %[[V2]], %[[M3]][3 : i32] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
 //       CHECK: llvm.return %[[M4]] : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
-  func @gpu_wmma_constant_op()  ->(!gpu.mma_matrix<16x16xf16, "COp">) {
+  builtin.func @gpu_wmma_constant_op()  ->(!gpu.mma_matrix<16x16xf16, "COp">) {
     %cst = constant 1.0 : f16
     %C = gpu.subgroup_mma_constant_matrix %cst : !gpu.mma_matrix<16x16xf16, "COp">
     return %C : !gpu.mma_matrix<16x16xf16, "COp">
   }
 }

mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir

 // RUN: mlir-opt %s -convert-gpu-to-rocdl -split-input-file | FileCheck %s
 // RUN: mlir-opt %s -convert-gpu-to-rocdl='index-bitwidth=32' -split-input-file | FileCheck --check-prefix=CHECK32 %s
 
 gpu.module @test_module {
   // CHECK-LABEL: func @gpu_index_ops()
   // CHECK32-LABEL: func @gpu_index_ops()
-  func @gpu_index_ops()
+  builtin.func @gpu_index_ops()
       -> (index, index, index, index, index, index,
           index, index, index, index, index, index) {
     // CHECK32-NOT: = llvm.sext %{{.*}} : i32 to i64
 
     // CHECK: rocdl.workitem.id.x : i32
     // CHECK: = llvm.sext %{{.*}} : i32 to i64
     %tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index)
     // CHECK: rocdl.workitem.id.y : i32
gpu.module @test_module {
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK-LABEL: func @gpu_index_comp
   // CHECK32-LABEL: func @gpu_index_comp
-  func @gpu_index_comp(%idx : index) -> index {
+  builtin.func @gpu_index_comp(%idx : index) -> index {
     // CHECK: = llvm.add %{{.*}}, %{{.*}} : i64
     // CHECK32: = llvm.add %{{.*}}, %{{.*}} : i32
     %0 = addi %idx, %idx : index
     // CHECK: llvm.return %{{.*}} : i64
     // CHECK32: llvm.return %{{.*}} : i32
     std.return %0 : index
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK-LABEL: func @gpu_sync()
-  func @gpu_sync() {
+  builtin.func @gpu_sync() {
     // CHECK: rocdl.barrier
     gpu.barrier
     std.return
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_fabs_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_fabs_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_fabs
-  func @gpu_fabs(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_fabs(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = std.absf %arg_f32 : f32
     // CHECK: llvm.call @__ocml_fabs_f32(%{{.*}}) : (f32) -> f32
     %result64 = std.absf %arg_f64 : f64
     // CHECK: llvm.call @__ocml_fabs_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_ceil_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_ceil_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_ceil
-  func @gpu_ceil(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_ceil(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = std.ceilf %arg_f32 : f32
     // CHECK: llvm.call @__ocml_ceil_f32(%{{.*}}) : (f32) -> f32
     %result64 = std.ceilf %arg_f64 : f64
     // CHECK: llvm.call @__ocml_ceil_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_floor_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_floor_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_floor
-  func @gpu_floor(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_floor(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = std.floorf %arg_f32 : f32
     // CHECK: llvm.call @__ocml_floor_f32(%{{.*}}) : (f32) -> f32
     %result64 = std.floorf %arg_f64 : f64
     // CHECK: llvm.call @__ocml_floor_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_cos_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_cos_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_cos
-  func @gpu_cos(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_cos(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.cos %arg_f32 : f32
     // CHECK: llvm.call @__ocml_cos_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.cos %arg_f64 : f64
     // CHECK: llvm.call @__ocml_cos_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_exp_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_exp_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_exp
-  func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %exp_f32 = math.exp %arg_f32 : f32
     // CHECK: llvm.call @__ocml_exp_f32(%{{.*}}) : (f32) -> f32
     %result32 = math.exp %exp_f32 : f32
     // CHECK: llvm.call @__ocml_exp_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.exp %arg_f64 : f64
     // CHECK: llvm.call @__ocml_exp_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_exp2_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_exp2_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_exp2
-  func @gpu_exp2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_exp2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %exp2_f32 = math.exp2 %arg_f32 : f32
     // CHECK: llvm.call @__ocml_exp2_f32(%{{.*}}) : (f32) -> f32
     %result32 = math.exp2 %exp2_f32 : f32
     // CHECK: llvm.call @__ocml_exp2_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.exp2 %arg_f64 : f64
     // CHECK: llvm.call @__ocml_exp2_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 // Test that we handled properly operation with SymbolTable other than module op
 gpu.module @test_module {
   "test.symbol_scope"() ({
     // CHECK: test.symbol_scope
     // CHECK: llvm.func @__ocml_exp_f32(f32) -> f32
     // CHECK: llvm.func @__ocml_exp_f64(f64) -> f64
     // CHECK-LABEL: func @gpu_exp
-    func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+    builtin.func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
       %exp_f32 = math.exp %arg_f32 : f32
       // CHECK: llvm.call @__ocml_exp_f32(%{{.*}}) : (f32) -> f32
       %result32 = math.exp %exp_f32 : f32
       // CHECK: llvm.call @__ocml_exp_f32(%{{.*}}) : (f32) -> f32
       %result64 = math.exp %arg_f64 : f64
       // CHECK: llvm.call @__ocml_exp_f64(%{{.*}}) : (f64) -> f64
       std.return %result32, %result64 : f32, f64
     }
     "test.finish" () : () -> ()
   }) : () -> ()
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_expm1_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_expm1_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_expm1
-  func @gpu_expm1(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_expm1(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %expm1_f32 = math.expm1 %arg_f32 : f32
     // CHECK: llvm.call @__ocml_expm1_f32(%{{.*}}) : (f32) -> f32
     %result32 = math.expm1 %expm1_f32 : f32
     // CHECK: llvm.call @__ocml_expm1_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.expm1 %arg_f64 : f64
     // CHECK: llvm.call @__ocml_expm1_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_log_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_log_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_log
-  func @gpu_log(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_log(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.log %arg_f32 : f32
     // CHECK: llvm.call @__ocml_log_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.log %arg_f64 : f64
     // CHECK: llvm.call @__ocml_log_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_log1p_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_log1p_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_log1p
-  func @gpu_log1p(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_log1p(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.log1p %arg_f32 : f32
     // CHECK: llvm.call @__ocml_log1p_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.log1p %arg_f64 : f64
     // CHECK: llvm.call @__ocml_log1p_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_log10_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_log10_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_log10
-  func @gpu_log10(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_log10(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.log10 %arg_f32 : f32
     // CHECK: llvm.call @__ocml_log10_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.log10 %arg_f64 : f64
     // CHECK: llvm.call @__ocml_log10_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_log2_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_log2_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_log2
-  func @gpu_log2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_log2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.log2 %arg_f32 : f32
     // CHECK: llvm.call @__ocml_log2_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.log2 %arg_f64 : f64
     // CHECK: llvm.call @__ocml_log2_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_rsqrt_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_rsqrt_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_rsqrt
-  func @gpu_rsqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
+  builtin.func @gpu_rsqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
       -> (f16, f32, f64) {
     %result16 = math.rsqrt %arg_f16 : f16
     // CHECK: llvm.fpext %{{.*}} : f16 to f32
     // CHECK-NEXT: llvm.call @__ocml_rsqrt_f32(%{{.*}}) : (f32) -> f32
     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
     %result32 = math.rsqrt %arg_f32 : f32
     // CHECK: llvm.call @__ocml_rsqrt_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.rsqrt %arg_f64 : f64
     // CHECK: llvm.call @__ocml_rsqrt_f64(%{{.*}}) : (f64) -> f64
     std.return %result16, %result32, %result64 : f16, f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_sqrt_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_sqrt_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_sqrt
-  func @gpu_sqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
+  builtin.func @gpu_sqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
       -> (f16, f32, f64) {
     %result16 = math.sqrt %arg_f16 : f16
     // CHECK: llvm.fpext %{{.*}} : f16 to f32
     // CHECK-NEXT: llvm.call @__ocml_sqrt_f32(%{{.*}}) : (f32) -> f32
     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
     %result32 = math.sqrt %arg_f32 : f32
     // CHECK: llvm.call @__ocml_sqrt_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.sqrt %arg_f64 : f64
     // CHECK: llvm.call @__ocml_sqrt_f64(%{{.*}}) : (f64) -> f64
     std.return %result16, %result32, %result64 : f16, f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_tanh_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_tanh_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_tanh
-  func @gpu_tanh(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_tanh(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.tanh %arg_f32 : f32
     // CHECK: llvm.call @__ocml_tanh_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.tanh %arg_f64 : f64
     // CHECK: llvm.call @__ocml_tanh_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_atan_f32(f32) -> f32
   // CHECK: llvm.func @__ocml_atan_f64(f64) -> f64
   // CHECK-LABEL: func @gpu_atan
-  func @gpu_atan(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_atan(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.atan %arg_f32 : f32
     // CHECK: llvm.call @__ocml_atan_f32(%{{.*}}) : (f32) -> f32
     %result64 = math.atan %arg_f64 : f64
     // CHECK: llvm.call @__ocml_atan_f64(%{{.*}}) : (f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_atan2_f32(f32, f32) -> f32
   // CHECK: llvm.func @__ocml_atan2_f64(f64, f64) -> f64
   // CHECK-LABEL: func @gpu_atan2
-  func @gpu_atan2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_atan2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.atan2 %arg_f32, %arg_f32 : f32
     // CHECK: llvm.call @__ocml_atan2_f32(%{{.*}}) : (f32, f32) -> f32
     %result64 = math.atan2 %arg_f64, %arg_f64 : f64
     // CHECK: llvm.call @__ocml_atan2_f64(%{{.*}}) : (f64, f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 
 // -----
 
 gpu.module @test_module {
   // CHECK: llvm.func @__ocml_pow_f32(f32, f32) -> f32
   // CHECK: llvm.func @__ocml_pow_f64(f64, f64) -> f64
   // CHECK-LABEL: func @gpu_pow
-  func @gpu_pow(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
+  builtin.func @gpu_pow(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
     %result32 = math.powf %arg_f32, %arg_f32 : f32
     // CHECK: llvm.call @__ocml_pow_f32(%{{.*}}, %{{.*}}) : (f32, f32) -> f32
     %result64 = math.powf %arg_f64, %arg_f64 : f64
     // CHECK: llvm.call @__ocml_pow_f64(%{{.*}}, %{{.*}}) : (f64, f64) -> f64
     std.return %result32, %result64 : f32, f64
   }
 }
 

mlir/test/Conversion/VectorToROCDL/vector-to-rocdl.mlir

 // RUN: mlir-opt %s -convert-vector-to-rocdl | FileCheck %s
 
 gpu.module @test_read{
-func @transfer_readx2(%A : memref<?xf32>, %base: index) -> vector<2xf32> {
+builtin.func @transfer_readx2(%A : memref<?xf32>, %base: index) -> vector<2xf32> {
   %f0 = constant 0.0: f32
   %f = vector.transfer_read %A[%base], %f0
       {permutation_map = affine_map<(d0) -> (d0)>} :
     memref<?xf32>, vector<2xf32>
   return %f: vector<2xf32>
 }
 // CHECK-LABEL: @transfer_readx2
 // CHECK: rocdl.buffer.load {{.*}} vector<2xf32>
 
-func @transfer_readx4(%A : memref<?xf32>, %base: index) -> vector<4xf32> {
+builtin.func @transfer_readx4(%A : memref<?xf32>, %base: index) -> vector<4xf32> {
   %f0 = constant 0.0: f32
   %f = vector.transfer_read %A[%base], %f0
       {permutation_map = affine_map<(d0) -> (d0)>} :
     memref<?xf32>, vector<4xf32>
   return %f: vector<4xf32>
 }
 // CHECK-LABEL: @transfer_readx4
 // CHECK: rocdl.buffer.load {{.*}} vector<4xf32>
 
-func @transfer_read_dwordConfig(%A : memref<?xf32>, %base: index) -> vector<4xf32> {
+builtin.func @transfer_read_dwordConfig(%A : memref<?xf32>, %base: index) -> vector<4xf32> {
   %f0 = constant 0.0: f32
   %f = vector.transfer_read %A[%base], %f0
       {permutation_map = affine_map<(d0) -> (d0)>} :
     memref<?xf32>, vector<4xf32>
   return %f: vector<4xf32>
 }
 // CHECK-LABEL: @transfer_read_dwordConfig
 // CHECK: %[[gep:.*]] = llvm.getelementptr {{.*}}
 // CHECK: [0, 0, -1, 159744]
 // CHECK: %[[i64:.*]] = llvm.ptrtoint %[[gep]]
 // CHECK: llvm.insertelement %[[i64]]
 }
 
 gpu.module @test_write{
-func @transfer_writex2(%A : memref<?xf32>, %B : vector<2xf32>, %base: index) {
+builtin.func @transfer_writex2(%A : memref<?xf32>, %B : vector<2xf32>, %base: index) {
   vector.transfer_write %B, %A[%base]
       {permutation_map = affine_map<(d0) -> (d0)>} :
     vector<2xf32>, memref<?xf32>
   return
 }
 // CHECK-LABEL: @transfer_writex2
 // CHECK: rocdl.buffer.store {{.*}} vector<2xf32>
 
-func @transfer_writex4(%A : memref<?xf32>, %B : vector<4xf32>, %base: index) {
+builtin.func @transfer_writex4(%A : memref<?xf32>, %B : vector<4xf32>, %base: index) {
   vector.transfer_write %B, %A[%base]
       {permutation_map = affine_map<(d0) -> (d0)>} :
     vector<4xf32>, memref<?xf32>
   return
 }
 // CHECK-LABEL: @transfer_writex4
 // CHECK: rocdl.buffer.store {{.*}} vector<4xf32>
 
-func @transfer_write_dwordConfig(%A : memref<?xf32>, %B : vector<2xf32>, %base: index) {
+builtin.func @transfer_write_dwordConfig(%A : memref<?xf32>, %B : vector<2xf32>, %base: index) {
   vector.transfer_write %B, %A[%base]
       {permutation_map = affine_map<(d0) -> (d0)>} :
     vector<2xf32>, memref<?xf32>
   return
 }
 // CHECK-LABEL: @transfer_write_dwordConfig
 // CHECK: %[[gep:.*]] = llvm.getelementptr {{.*}}
 // CHECK: [0, 0, -1, 159744]
 // CHECK: %[[i64:.*]] = llvm.ptrtoint %[[gep]]
 // CHECK: llvm.insertelement %[[i64]]
 }

mlir/test/Dialect/Builtin/canonicalize.mlir

 // RUN: mlir-opt %s -canonicalize | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // UnrealizedConversionCastOp
 //===----------------------------------------------------------------------===//
 
 // Test folding conversion casts feeding into other casts.
 // CHECK-LABEL: func @multiple_conversion_casts
 // CHECK-SAME: %[[ARG0:.*]]: i32, %[[ARG1:.*]]:
 func @multiple_conversion_casts(%arg0: i32, %arg1: i32) -> (i32, i32) {
   // CHECK-NOT: unrealized_conversion_cast
   // CHECK: return %[[ARG0]], %[[ARG1]]
-  %inputs:2 = unrealized_conversion_cast %arg0, %arg1 : i32, i32 to i64, i64
-  %outputs:2 = unrealized_conversion_cast %inputs#0, %inputs#1 : i64, i64 to i32, i32
+  %inputs:2 = builtin.unrealized_conversion_cast %arg0, %arg1 : i32, i32 to i64, i64
+  %outputs:2 = builtin.unrealized_conversion_cast %inputs#0, %inputs#1 : i64, i64 to i32, i32
   return %outputs#0, %outputs#1 : i32, i32
 }
 
 // CHECK-LABEL: func @multiple_conversion_casts
 func @multiple_conversion_casts_failure(%arg0: i32, %arg1: i32, %arg2: i64) -> (i32, i32) {
   // CHECK: unrealized_conversion_cast
   // CHECK: unrealized_conversion_cast
-  %inputs:2 = unrealized_conversion_cast %arg0, %arg1 : i32, i32 to i64, i64
-  %outputs:2 = unrealized_conversion_cast %arg2, %inputs#1 : i64, i64 to i32, i32
+  %inputs:2 = builtin.unrealized_conversion_cast %arg0, %arg1 : i32, i32 to i64, i64
+  %outputs:2 = builtin.unrealized_conversion_cast %arg2, %inputs#1 : i64, i64 to i32, i32
   return %outputs#0, %outputs#1 : i32, i32
 }

mlir/test/Dialect/Linalg/drop-unit-extent-dims.mlir

outs(%shape : memref<?x1x?x1x?xf32>) {
        }
   return %shape : memref<?x1x?x1x?xf32>
 }
 
 // CHECK:     #[[MAP0:.*]] = affine_map<(d0, d1, d2)[s0] -> (d0 * s0 + d1 + d2)>
 // CHECK:     #[[MAP1:.*]] = affine_map<(d0, d1, d2) -> (d0, 0, d2)>
 // CHECK:     #[[MAP2:.*]] = affine_map<(d0, d1, d2) -> ()>
 // CHECK:     #[[MAP3:.*]] = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
-// CHECK:     builtin.func @input_stays_same(
+// CHECK:     func @input_stays_same(
 // CHECK-SAME:  %[[ARG0:.*]]: memref<?x1x?xf32, #[[MAP0]]>,
 // CHECK-SAME:  %[[ARG1:.*]]: f32, %[[ARG2:.*]]: memref<?x1x?x1x?xf32>)
 // CHECK-SAME   -> memref<?x1x?x1x?xf32> {
 // CHECK:      %[[OUT:.*]] = memref.collapse_shape %[[ARG2]] {{\[}}[0, 1], [2, 3], [4]]
 // CHECK-SAME:   : memref<?x1x?x1x?xf32> into memref<?x?x?xf32>
 // CHECK:      linalg.generic
 // CHECK-SAME:   {indexing_maps = [#[[MAP1]], #[[MAP2]], #[[MAP3]]],
 // CHECK-SAME:   iterator_types = ["parallel", "parallel", "parallel"]}
 // CHECK-SAME:   ins(%[[ARG0]], %[[ARG1]] : memref<?x1x?xf32, #[[MAP0]]>, f32)
 // CHECK-SAME:   outs(%[[OUT]] : memref<?x?x?xf32>) {
 // CHECK:      ^bb0(%{{.*}}: f32, %[[ARG:.*]]: f32, %{{.*}}: f32): // no predecessors
 // CHECK:       linalg.yield %[[ARG]] : f32
 // CHECK:      }
 // CHECK:      return %[[ARG2]] : memref<?x1x?x1x?xf32>

mlir/test/Dialect/Linalg/reshape_control_fusion.mlir

%1 = linalg.generic {
       ^bb0(%arg2 : f32, %arg3:f32, %arg4 : f32):
         %2 = addf %arg2, %arg3 : f32
         linalg.yield %2 : f32
       } -> tensor<?x?xf32>
   return %1 : tensor<?x?xf32>
 }
 //  CHECK-DAG: #[[MAP0:.+]] = affine_map<(d0, d1) -> (d0, d1)>
 //  CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1) -> (d1)>
-//      CHECK: builtin.func @control_producer_reshape_fusion
+//      CHECK: func @control_producer_reshape_fusion
 // CHECK-SAME:   %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?x?xf32>
 // CHECK-SAME:   %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?xf32>
 //  CHECK-DAG:   %[[C0:.+]] = constant 0 : index
 //  CHECK-DAG:   %[[C1:.+]] = constant 1 : index
 //      CHECK:   %[[RESHAPE:.+]] = linalg.tensor_collapse_shape %[[ARG0]]
 // CHECK-SAME:       {{\[}}[0, 1], [2]{{\]}} : tensor<?x?x?xf32> into tensor<?x?xf32>
 //      CHECK:   %[[RESULT:.+]] = linalg.generic
 // CHECK-SAME:       indexing_maps = [#[[MAP0]], #[[MAP1]], #[[MAP0]]]
%fill = linalg.generic {
         linalg.yield %cst : f32
       } -> tensor<?x?xf32>
   %0 = linalg.tensor_expand_shape %fill [[0, 1], [2]] : tensor<?x?xf32> into tensor<1x?x?xf32>
   %1 = linalg.batch_matmul ins(%arg0, %arg1 : tensor<1x?x?xf32>, tensor<1x?x?xf32>)
       outs(%0 : tensor<1x?x?xf32>) -> tensor<1x?x?xf32>
   return %1 : tensor<1x?x?xf32>
 }
 //  CHECK-DAG: #[[MAP:.+]] = affine_map<(d0, d1, d2) -> (d0, d1, d2)
-//      CHECK: builtin.func @control_consumer_reshape_fusion
+//      CHECK: func @control_consumer_reshape_fusion
 //      CHECK:   %[[FILL:.+]] = linalg.generic
 // CHECK-SAME:       indexing_maps = [#[[MAP]]]
 // CHECK-SAME:       outs(%{{.+}} : tensor<1x?x?xf32>)
 //      CHECK:   linalg.batch_matmul
 // CHECK-SAME:       outs(%[[FILL]] : tensor<1x?x?xf32>)

mlir/test/Dialect/Shape/invalid.mlir

 
 // Test using an unsupported shape.lib attribute type.
 
 // expected-error@+1 {{only SymbolRefAttr allowed in shape.lib attribute array}}
 module attributes {shape.lib = [@shape_lib, "shape_lib"]} {
 
 shape.function_library @shape_lib {
   // Test shape function that returns the shape of input arg as result shape.
-  func @same_result_shape(%arg: !shape.value_shape) -> !shape.shape {
+  builtin.func @same_result_shape(%arg: !shape.value_shape) -> !shape.shape {
     %0 = shape.shape_of %arg : !shape.value_shape -> !shape.shape
     return %0 : !shape.shape
   }
 } mapping {
   test.same_operand_result_type = @same_result_shape
 }
 
 }
 
 // -----
 
 // Test that duplicate op to shape function mappings are flagged, this uses
 // the same library twice for easy overlap.
 
 // expected-error@+1 {{only one op to shape mapping allowed}}
 module attributes {shape.lib = [@shape_lib, @shape_lib]} {
 
 shape.function_library @shape_lib {
   // Test shape function that returns the shape of input arg as result shape.
-  func @same_result_shape(%arg: !shape.value_shape) -> !shape.shape {
+  builtin.func @same_result_shape(%arg: !shape.value_shape) -> !shape.shape {
     %0 = shape.shape_of %arg : !shape.value_shape -> !shape.shape
     return %0 : !shape.shape
   }
 } mapping {
   test.same_operand_result_type = @same_result_shape
 }
 
 }
 
 // -----
 
 // Test that duplicate op to shape function mappings are flagged (this is
 // more an invariant of using the dictionary attribute here than anything
 // specific to function library op).
 
 module attributes {shape.lib = [@shape_lib]} {
 
 shape.function_library @shape_lib {
   // Test shape function that returns the shape of input arg as result shape.
-  func @same_result_shape(%arg: !shape.value_shape) -> !shape.shape {
+  builtin.func @same_result_shape(%arg: !shape.value_shape) -> !shape.shape {
     %0 = shape.shape_of %arg : !shape.value_shape -> !shape.shape
     return %0 : !shape.shape
   }
 } mapping {
   // expected-error @+2 {{duplicate key}}
   test.same_operand_result_type = @same_result_shape,
   test.same_operand_result_type = @same_result_shape
 }

mlir/test/Dialect/SparseTensor/sparse_perm.mlir

 // NOTE: Assertions have been autogenerated by utils/generate-test-checks.py
 // RUN: mlir-opt %s -sparsification | FileCheck %s
 
 #X = #sparse_tensor.encoding<{
  dimLevelType = [ "dense", "dense", "dense" ],
  dimOrdering = affine_map<(i,j,k) -> (k,i,j)>
 }>
 
 #trait = {
   indexing_maps = [
     affine_map<(i,j,k) -> (k,i,j)>,  // A (in)
     affine_map<(i,j,k) -> (i,j,k)>   // X (out)
   ],
   iterator_types = ["parallel", "parallel", "parallel"]
 }
 
-// CHECK-LABEL:   builtin.func @sparse_static_dims(
+// CHECK-LABEL:   func @sparse_static_dims(
 // CHECK-SAME:                          %[[VAL_0:.*]]: tensor<10x20x30xf32, #sparse_tensor.encoding<{{{.*}}}>>,
 // CHECK-SAME:                          %[[VAL_1:.*]]: tensor<20x30x10xf32>) -> tensor<20x30x10xf32> {
 // CHECK:           %[[VAL_2:.*]] = constant 20 : index
 // CHECK:           %[[VAL_3:.*]] = constant 30 : index
 // CHECK:           %[[VAL_4:.*]] = constant 10 : index
 // CHECK:           %[[VAL_5:.*]] = constant 0 : index
 // CHECK:           %[[VAL_6:.*]] = constant 1 : index
 // CHECK:           %[[VAL_7:.*]] = sparse_tensor.values %[[VAL_0]] : tensor<10x20x30xf32, #sparse_tensor.encoding<{{{.*}}}>>
%0 = linalg.generic #trait
     ins(%arga: tensor<10x20x30xf32, #X>)
     outs(%argx: tensor<20x30x10xf32>) {
       ^bb(%a : f32, %x: f32):
         linalg.yield %a : f32
   } -> tensor<20x30x10xf32>
   return %0 : tensor<20x30x10xf32>
 }
 
-// CHECK-LABEL:   builtin.func @sparse_dynamic_dims(
+// CHECK-LABEL:   func @sparse_dynamic_dims(
 // CHECK-SAME:                          %[[VAL_0:.*]]: tensor<?x?x?xf32, #sparse_tensor.encoding<{{{.*}}}>>,
 // CHECK-SAME:                          %[[VAL_1:.*]]: tensor<?x?x?xf32>) -> tensor<?x?x?xf32> {
 // CHECK:           %[[VAL_2:.*]] = constant 2 : index
 // CHECK:           %[[VAL_3:.*]] = constant 0 : index
 // CHECK:           %[[VAL_4:.*]] = constant 1 : index
 // CHECK:           %[[VAL_5:.*]] = sparse_tensor.values %[[VAL_0]] : tensor<?x?x?xf32, #sparse_tensor.encoding<{{{.*}}}>>
 // CHECK:           %[[VAL_6:.*]] = tensor.dim %[[VAL_1]], %[[VAL_3]] : tensor<?x?x?xf32>
 // CHECK:           %[[VAL_7:.*]] = tensor.dim %[[VAL_1]], %[[VAL_4]] : tensor<?x?x?xf32>

mlir/test/Dialect/SparseTensor/sparse_perm_lower.mlir

 #trait = {
   indexing_maps = [
     affine_map<(i,j,k) -> (k,i,j)>,  // A (in)
     affine_map<(i,j,k) -> ()>        // X (out)
   ],
   iterator_types = ["reduction", "reduction", "reduction"]
 }
 
-// CHECK-HIR-LABEL:   builtin.func @sparse_dynamic_dims(
+// CHECK-HIR-LABEL:   func @sparse_dynamic_dims(
 // CHECK-HIR-SAME:                                      %[[VAL_0:.*]]: tensor<?x?x?xf32,  #sparse_tensor.encoding<{{{.*}}}>>,
 // CHECK-HIR-SAME:                                      %[[VAL_1:.*]]: tensor<f32>) -> tensor<f32> {
 // CHECK-HIR-DAG:       %[[C0:.*]] = constant 0 : index
 // CHECK-HIR-DAG:       %[[C1:.*]] = constant 1 : index
 // CHECK-HIR-DAG:       %[[C2:.*]] = constant 2 : index
 // CHECK-HIR:           %[[VAL_5:.*]] = tensor.dim %[[VAL_0]], %[[C2]] : tensor<?x?x?xf32,  #sparse_tensor.encoding<{{{.*}}}>>
 // CHECK-HIR:           %[[VAL_6:.*]] = tensor.dim %[[VAL_0]], %[[C0]] : tensor<?x?x?xf32,  #sparse_tensor.encoding<{{{.*}}}>>
 // CHECK-HIR:           %[[VAL_7:.*]] = tensor.dim %[[VAL_0]], %[[C1]] : tensor<?x?x?xf32,  #sparse_tensor.encoding<{{{.*}}}>>

 // CHECK-HIR:               }
 // CHECK-HIR:               memref.store %[[VAL_23:.*]], %[[VAL_10]][] : memref<f32>
 // CHECK-HIR:             }
 // CHECK-HIR:           }
 // CHECK-HIR:           %[[VAL_24:.*]] = memref.tensor_load %[[VAL_10]] : memref<f32>
 // CHECK-HIR:           return %[[VAL_24]] : tensor<f32>
 // CHECK-HIR:         }
 //
-// CHECK-MIR-LABEL:   builtin.func @sparse_dynamic_dims(
+// CHECK-MIR-LABEL:   func @sparse_dynamic_dims(
 // CHECK-MIR-SAME:                                      %[[VAL_0:.*]]: !llvm.ptr<i8>,
 // CHECK-MIR-SAME:                                      %[[VAL_1:.*]]: tensor<f32>) -> tensor<f32> {
 // CHECK-MIR-DAG:       %[[C0:.*]] = constant 0 : index
 // CHECK-MIR-DAG:       %[[C1:.*]] = constant 1 : index
 // CHECK-MIR-DAG:       %[[C2:.*]] = constant 2 : index
 // CHECK-MIR:           %[[VAL_5:.*]] = call @sparseDimSize(%[[VAL_0]], %[[C0]]) : (!llvm.ptr<i8>, index) -> index
 // CHECK-MIR:           %[[VAL_6:.*]] = call @sparseDimSize(%[[VAL_0]], %[[C1]]) : (!llvm.ptr<i8>, index) -> index
 // CHECK-MIR:           %[[VAL_7:.*]] = call @sparseDimSize(%[[VAL_0]], %[[C2]]) : (!llvm.ptr<i8>, index) -> index

mlir/test/IR/invalid-func-op.mlir

 // RUN: mlir-opt %s -split-input-file -verify-diagnostics
 
 // -----
 
 func @func_op() {
   // expected-error@+1 {{expected valid '@'-identifier for symbol name}}
-  func missingsigil() -> (i1, index, f32)
+  builtin.func missingsigil() -> (i1, index, f32)
   return
 }
 
 // -----
 
 func @func_op() {
   // expected-error@+1 {{expected type instead of SSA identifier}}
-  func @mixed_named_arguments(f32, %a : i32) {
+  builtin.func @mixed_named_arguments(f32, %a : i32) {
     return
   }
   return
 }
 
 // -----
 
 func @func_op() {
   // expected-error@+1 {{expected SSA identifier}}
-  func @mixed_named_arguments(%a : i32, f32) -> () {
+  builtin.func @mixed_named_arguments(%a : i32, f32) -> () {
     return
   }
   return
 }
 
 // -----
 
 func @func_op() {
   // expected-error@+1 {{entry block must have 1 arguments to match function signature}}
-  func @mixed_named_arguments(f32) {
+  builtin.func @mixed_named_arguments(f32) {
   ^entry:
     return
   }
   return
 }
 
 // -----
 
 func @func_op() {
   // expected-error@+1 {{type of entry block argument #0('i32') must match the type of the corresponding argument in function signature('f32')}}
-  func @mixed_named_arguments(f32) {
+  builtin.func @mixed_named_arguments(f32) {
   ^entry(%arg : i32):
     return
   }
   return
 }
 
 // -----
 

mlir/test/IR/invalid-module-op.mlir

 // RUN: mlir-opt %s -split-input-file -verify-diagnostics
 
 // -----
 
 func @module_op() {
   // expected-error@+1 {{Operations with a 'SymbolTable' must have exactly one block}}
-  module {
+  builtin.module {
   ^bb1:
     "test.dummy"() : () -> ()
   ^bb2:
     "test.dummy"() : () -> ()
   }
   return
 }
 
 // -----
 
 func @module_op() {
   // expected-error@+1 {{region should have no arguments}}
-  module {
+  builtin.module {
   ^bb1(%arg: i32):
   }
   return
 }
 
 // -----
 
 // expected-error@+1 {{can only contain attributes with dialect-prefixed names}}
 module attributes {attr} {
 }
 

mlir/test/IR/invalid-ops.mlir

^bb0:
   %x = "affine.apply" (%i, %j) {map = affine_map<(d0, d1) -> ((d0 + 1), (d1 + 2))>} : (index,index) -> (index) //  expected-error {{'affine.apply' op mapping must produce one value}}
   return
 }
 
 // -----
 
 func @unknown_custom_op() {
 ^bb0:
-  %i = crazyThing() {value = 0} : () -> index  // expected-error {{custom op 'crazyThing' is unknown}}
+  %i = test.crazyThing() {value = 0} : () -> index  // expected-error {{custom op 'test.crazyThing' is unknown}}
   return
 }
 
 // -----
 
 func @unknown_std_op() {
   // expected-error@+1 {{unregistered operation 'std.foo_bar_op' found in dialect ('std') that does not allow unknown operations}}
   %0 = "std.foo_bar_op"() : () -> index

mlir/test/IR/invalid.mlir

 
 func @for_negative_stride() {
   affine.for %i = 1 to 10 step -1
 }        // expected-error@-1 {{expected step to be representable as a positive signed integer}}
 
 // -----
 
 func @non_operation() {
-  asd   // expected-error {{custom op 'asd' is unknown}}
+  test.asd   // expected-error {{custom op 'test.asd' is unknown}}
 }
 
 // -----
 
 func @invalid_if_conditional2() {
   affine.for %i = 1 to 10 {
     affine.if affine_set<(i)[N] : (i >= )>  // expected-error {{expected '== 0' or '>= 0' at end of affine constraint}}
   }

mlir/test/IR/parser.mlir

func @pretty_names() {
 
   %q:3, %r = test.string_attr_pretty_name
   // CHECK: %q, %q_1, %q_2, %r = test.string_attr_pretty_name attributes {names = ["q", "q", "q", "r"]}
 
   // CHECK: return
   return
 }
 
+
+// This tests the behavior of "default dialect":
+// operations like `test.default_dialect` can define a default dialect
+// used in nested region.
+// CHECK-LABEL: func @default_dialect
+func @default_dialect() {
+  test.default_dialect {
+    // The test dialect is the default in this region, the following two
+    // operations are parsed identically.
+    // CHECK-NOT: test.parse_integer_literal
+    parse_integer_literal : 5
+    // CHECK: parse_integer_literal : 6
+    test.parse_integer_literal : 6
+    // Verify that only an op prefix is stripped, not an attribute value for
+    // example.
+    // CHECK:  "test.op_with_attr"() {test.attr = "test.value"} : () -> ()
+    "test.op_with_attr"() {test.attr = "test.value"} : () -> ()
+    "test.terminator"() : ()->()
+  }
+  return
+}
+
 // CHECK-LABEL: func @unreachable_dominance_violation_ok
 func @unreachable_dominance_violation_ok() -> i1 {
 // CHECK:   [[VAL:%.*]] = constant false
 // CHECK:   return [[VAL]] : i1
 // CHECK: ^bb1:   // no predecessors
 // CHECK:   [[VAL2:%.*]]:3 = "bar"([[VAL3:%.*]]) : (i64) -> (i1, i1, i1)
 // CHECK:   br ^bb3
 // CHECK: ^bb2:   // pred: ^bb2

mlir/test/IR/traits.mlir

 }
 
 // -----
 
 // Ensure that SSACFG regions of operations in GRAPH regions are
 // checked for dominance
 func @illegalInsideDominanceFreeScope() -> () {
   test.graph_region {
-    func @test() -> i1 {
+    builtin.func @test() -> i1 {
     ^bb1:
       // expected-error @+1 {{operand #0 does not dominate this use}}
       %2:3 = "bar"(%1) : (i64) -> (i1,i1,i1)
       // expected-note @+1 {{operand defined here}}
      %1 = "baz"(%2#0) : (i1) -> (i64)
       return %2#1 : i1
     }
     "terminator"() : () -> ()
   }
   return
 }
 
 // -----
 
 // Ensure that SSACFG regions of operations in GRAPH regions are
 // checked for dominance
 func @illegalCDFGInsideDominanceFreeScope() -> () {
   test.graph_region {
-    func @test() -> i1 {
+    builtin.func @test() -> i1 {
     ^bb1:
       // expected-error @+1 {{operand #0 does not dominate this use}}
       %2:3 = "bar"(%1) : (i64) -> (i1,i1,i1)
       br ^bb4
     ^bb2:
       br ^bb2
     ^bb4:
       %1 = "foo"() : ()->i64   // expected-note {{operand defined here}}

mlir/test/Transforms/canonicalize-dce.mlir

 
 // Test case: Recursively DCE into enclosed regions.
 
 // CHECK:      func @f(%arg0: f32)
 // CHECK-NEXT:   func @g(%arg1: f32)
 // CHECK-NEXT:     return
 
 func @f(%arg0: f32) {
-  func @g(%arg1: f32) {
+  builtin.func @g(%arg1: f32) {
     %0 = "std.addf"(%arg1, %arg1) : (f32, f32) -> f32
     return
   }
   return
 }
 
 // -----
 

mlir/test/Transforms/canonicalize.mlir

func @write_only_alloca_fold(%v: f32) {
   %a = memref.alloca(%c4) : memref<?xf32>
   memref.store %v, %a[%c0] : memref<?xf32>
   return
 }
 
 // CHECK-LABEL: func @dead_block_elim
 func @dead_block_elim() {
   // CHECK-NOT: ^bb
-  func @nested() {
+  builtin.func @nested() {
     return
 
   ^bb1:
     return
   }
   return
 
 ^bb1:

mlir/test/Transforms/constant-fold.mlir

 }
 
 // -----
 
 // CHECK-LABEL: func @nested_isolated_region
 func @nested_isolated_region() {
   // CHECK-NEXT: func @isolated_op
   // CHECK-NEXT: constant 2
-  func @isolated_op() {
+  builtin.func @isolated_op() {
     %0 = constant 1 : i32
     %2 = addi %0, %0 : i32
     "foo.yield"(%2) : (i32) -> ()
   }
 
   // CHECK: "foo.unknown_region"
   // CHECK-NEXT: constant 2
   "foo.unknown_region"() ({

mlir/test/Transforms/cse.mlir

 /// This test checks that nested regions that are isolated from above are
 /// properly handled.
 // CHECK-LABEL: @nested_isolated
 func @nested_isolated() -> i32 {
   // CHECK-NEXT: constant 1
   %0 = constant 1 : i32
 
   // CHECK-NEXT: @nested_func
-  func @nested_func() {
+  builtin.func @nested_func() {
     // CHECK-NEXT: constant 1
     %foo = constant 1 : i32
     "foo.yield"(%foo) : (i32) -> ()
   }
 
   // CHECK: "foo.region"
   "foo.region"() ({
     // CHECK-NEXT: constant 1

mlir/test/Transforms/test-legalizer-full.mlir

func @replace_non_root_illegal_op() {
   "test.return"() : () -> ()
 }
 
 // -----
 
 // Test that children of recursively legal operations are ignored.
 func @recursively_legal_invalid_op() {
   /// Operation that is statically legal.
-  module attributes {test.recursively_legal} {
+  builtin.module attributes {test.recursively_legal} {
     %ignored = "test.illegal_op_f"() : () -> (i32)
   }
   /// Operation that is dynamically legal, i.e. the function has a pattern
   /// applied to legalize the argument type before it becomes recursively legal.
-  func @dynamic_func(%arg: i64) attributes {test.recursively_legal} {
+  builtin.func @dynamic_func(%arg: i64) attributes {test.recursively_legal} {
     %ignored = "test.illegal_op_f"() : () -> (i32)
     "test.return"() : () -> ()
   }
 
   "test.return"() : () -> ()
 }
 
 // -----

mlir/test/lib/Dialect/Test/TestOps.td

let results = (outs
     Variadic<I32>:$d
   );
 }
 
 // This is used to test encoding of a string attribute into an SSA name of a
 // pretty printed value name.
 def StringAttrPrettyNameOp
  : TEST_Op<"string_attr_pretty_name",
-           [DeclareOpInterfaceMethods<OpAsmOpInterface>]> {
+           [DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>]> {
   let arguments = (ins StrArrayAttr:$names);
   let results = (outs Variadic<I32>:$r);
 
   let printer = [{ return ::print(p, *this); }];
   let parser = [{ return ::parse$cppClass(parser, result); }];
 }
 
+// This is used to test the OpAsmOpInterface::getDefaultDialect() feature:
+// operations nested in a region under this op will drop the "test." dialect
+// prefix.
+def DefaultDialectOp : TEST_Op<"default_dialect", [OpAsmOpInterface]> {
+ let regions = (region AnyRegion:$body);
+  let extraClassDeclaration = [{
+    static ::llvm::StringRef getDefaultDialect() {
+      return "test";
+    }
+    void getAsmResultNames(::llvm::function_ref<void(::mlir::Value, ::llvm::StringRef)> setNameFn) {}
+  }];
+  let assemblyFormat = "regions attr-dict-with-keyword";
+}
+
 //===----------------------------------------------------------------------===//
 // Test Locations
 //===----------------------------------------------------------------------===//
 
 def TestLocationSrcOp : TEST_Op<"loc_src"> {
   let arguments = (ins I32:$input);
   let results = (outs I32:$output);
 }