[mlir][spirv] Introduce spv.function

[antiagainst]

Summary

Thus far we have been using builtin func op to model SPIR-V functions.
It was because builtin func op used to have special treatment in
various parts of the core codebase (e.g., pass pipelines, etc.) and
it's easy to bootstrap the development of the SPIR-V dialect. But
nowadays with general op concepts and region support we don't have
such limitations and it's time to tighten the SPIR-V dialect for
completeness.

This commits introduces a spv.function op to properly model SPIR-V
functions. Compared to builtin func op, it can provide the following
benefits:

• We can control the full op so we can integrate SPIR-V information bits (e.g., function control) in a more integrated way and define our own assembly form and enforcing better verification.
• We can have a better dialect and library boundary. At the current moment only functions are modelled with an external op. With this change, all ops modelling SPIR-V concpets will be spv.* ops and registered to the SPIR-V dialect.
• We don't need to special-case func op anymore when creating ConversionTarget declaring SPIR-V dialect as legal. This is quite important given we'll see more and more conversions in the future.

The operation is explicitly named differently than func op to
avoid potential confusion and symbol collision.

[rriddle]

I would actually prefer that we call this FuncOp('func'). It is consistent with all of the other functions that have been defined, and seems a little weird to deviate. For example, spirv::ModuleOp/ConstantOp are also named the same as other common ops.

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h

49

Could we auto-generate these in OpDefGen?

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

79

if (auto funcOp = ...)

return funcOp.getType();

return FunctionType();

?

1724

Drop trivial brace.

[mehdi_amini]

In D74226#1865003, @rriddle wrote:

I would actually prefer that we call this FuncOp('func'). It is consistent with all of the other functions that have been defined, and seems a little weird to deviate.

+1

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

277

The 4 methods above all seem like they could be added to FunctionLike trait?

287

getType().getNumResults(); is shorter?

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

1687

Seems like this should be checked in verifyType()?

I'd expect verifyBody to actually find the return ops and check for the matching, as well as other possible SPIRV specific properties.

[mravishankar]

THanks for taking this one now. Its already touching so much code, if we left it for longer it would be more effort to use spv.function

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

251

I think we can leave hasOpcode = 1. The autogenSerialization = 0 should avoid the serialization method to from being autogenerated.

mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRVPass.cpp

40–41

I dont think we need this. This is only for testing. So not adding this, but leaving the conversion as applyPartialConversion is fine.

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

2525

I dont think we should do this. spv.return must stay in spv.func. If for testing we can use func and return and the body can have spv.* operations. We just need to not do applyFullConversion and it should be fine.

mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp

1645

If you set hasOpCode = 1 in the SPIRVBase.td this should be generated by ODS automatically.

[antiagainst]

Address comments

[antiagainst]

I would actually prefer that we call this FuncOp('func'). It is consistent with all of the other functions that have been defined, and seems a little weird to deviate. For example, spirv::ModuleOp/ConstantOp are also named the same as other common ops.

The purpose was trying to avoid potential confusions between std ones but you are right this is already happening for other SPIR-V ops. So consistency wins. Done.

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h

49

Yup we can. But do we want to do that for each op? With so many ops having this specialization, it can increase C++ compilers' burden quite a bit. I'd defer that to when there are more ops wanting this.

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

251

We cannot here because this is not a direct mapping.

277

Other function-like ops (like gpu.func and llvm.func) might want slightly different behavior for some of the ops but yes this can be moved to FunctionLike to be a default impl. Done.

mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRVPass.cpp

40–41

It's tricky here. We need to pull in func to spv.func conversion so that we can materialize function signature conversion, which is needed for load/store op conversion in subview tests. Once that pattern is pulled in, it might cause that some test IR is in an invalid state. For example, mulf %arg, %arg: tensor<4xf32> won't be converted because we don't have a pattern to convert binary ops that require type conversion. But the %arg will be converted anyway because of func to spv.func change. So we will see a validation failure on mulf. I was trying to address this by turning the pass to do full conversion here given this pass is mainly for testing so we don't care that much. I think another way to address this is to disable the test for mulf %arg, %arg: tensor<4xf32>. It's actually just a negative test that checks we are missing a pattern. Done.

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

1687

Good points. Moved this check to verifyBody and bumped checks for return ops here.

2525

Bumped the check to spv.function so we don't need special-casing the enclosing function types here. But yes it's allowing return ops to be used at more places, which is more friendly to interop between different dialects. If ops from other dialects would like to use SPIR-V return ops then it's their responsibility to make sure the semantics match with their cases. This is consistent with how we handle ops within SPIR-V dialect: we allow non-region ops to be mixable with other dialects but we enforce checks in region ops (like spv.module, spv.func) to verify SPIR-V semantics at the scope of those ops.

More reasons of allowing return ops to be interop with other dialect ops: ideally I'd like to enforce the spv.func op to verify its region to only contain SPIR-V ops, and I'd like to make the func to spv.func op conversion to only kick in when a func op's region already only contains SPIR-V ops. That means, we need to allow other ops that can be placed inside functions to be converted first. If instead we require SPIR-V return ops to only work with spv.func, that causes a circular dependency... Theoretically it can be solved by throwing all patterns in the same pass maybe, but I'm not sure we want to go that far right now. Besides, func and std.return actually are in different dialects and likely we will have pass boundaries between their patterns anyway.

mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp

1645

We cannot here because this is not a direct mapping: the op is not named as spv.Function; it was named spv.function and now is spv.func.

[mravishankar]

Thanks Lei! Accepting conditioned on addressing other comments.

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

251

Does it not map directly to OpFunction . If you name the op spv.Function , then the ODS will generate the dispatch for the serialization and deserialization. You just need to specialize the implementation of processOp in the (de)serializer for this to work. (See comment here : https://github.com/llvm/llvm-project/blob/master/mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp#L1664)
Anyway, this would be nice to have, is not necessary

mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp

1645

Why not name it spv.Function ? Curious. Wont make it a requirement. THis is fine too.

[rriddle]

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h

48

///

49

This needs to be formatted.

mlir/include/mlir/IR/FunctionSupport.h

148

None of these are valid if the type of the function isn't the builtin FunctionType, i.e. LLVMFuncOp doesn't work with any of these methods.

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

1693

You can return the errors directly. failure implicitly converts to WalkResult::interrupt. You just need to set the return type on the lambda to WalkResult.

[rriddle]

mlir/include/mlir/IR/FunctionSupport.h

148

If you want to add these, the things that rely on FunctionType should be moved to a separate section and there should be disclaimers/comments on the top-level comment of FunctionLike denoting that these methods need to be hidden/overridden if the underlying type isn't FunctionType. You should also make sure that these are all hidden for LLVMFuncOp.

[rriddle]

mlir/include/mlir/IR/FunctionSupport.h

148

An alternative that may work is to static assert that the type returned by std::declval<ConcreteType>().getType() is equal to FunctionType inside of each of those methods.

[rriddle]

LGTM after comments are resolved.

[antiagainst]

Address comments

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h

49

This is actually the result after running format fix internally. Maybe because clang-format version differences.

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

251

OpFunction is just one instruction composing the whole function; others include OpFunctionParameter, OpFunctionEnd. spv.func here actually includes all of them. We have been following the convention to use lower case names for such cases.

mlir/include/mlir/IR/FunctionSupport.h

148

Yup. But won't this still be good defaults for other func ops (std, gpu, spv)? llvm side one can hide these methods. (And that's what happening per the current impl.) Let me know if this is okay, I can always put these back and copy on the SPIR-V side.

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

1693

Oh, that's nice! Thanks for the suggestion! Changed.

mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp

1645

Explained at another place.

[antiagainst]

mlir/include/mlir/IR/FunctionSupport.h

148

Oh, replied before seeing your new comments. Sure, I'll address these. :)

[antiagainst]

Add comments on FunctionLike

[antiagainst]

mlir/include/mlir/IR/FunctionSupport.h

148

Added comments to explain for now. Tried the following

static_assert(
    std::is_same<FunctionType,
                 decltype(std::declval<ConcreteType>().getType())>::value,
    "must be hidden by the concrete type");

But didn't seem to guard anything. (Removing the LLVM overrides still can compile successfully.)

[mehdi_amini]

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h

49

Google internal does not honor .clang-format files: source of truth is the bot at the moment. Please use git clang-format

[antiagainst]

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h

49

Oh that's good to know! Thanks Mehdi! I formatted this.

mlir/include/mlir/Dialect/GPU/GPUOps.td

let extraClassDeclaration = [{
 
     /// Returns `true` if the GPU function defined by this Op is a kernel, i.e.
     /// it is intended to be launched from host.
     bool isKernel() {
       return getAttrOfType<UnitAttr>(GPUDialect::getKernelFuncAttrName()) !=
              nullptr;
     }
 
-    /// Returns the type of the function this Op defines.
-    FunctionType getType() {
-      return getTypeAttr().getValue().cast<FunctionType>();
-    }
-
     /// Change the type of this function in place. This is an extremely
     /// dangerous operation and it is up to the caller to ensure that this is
     /// legal for this function, and to restore invariants:
     ///  - the entry block args must be updated to match the function params.
     ///  - the argument/result attributes may need an update: if the new type
     ///  has less parameters we drop the extra attributes, if there are more
     ///  parameters they won't have any attributes.
     // TODO(b/146349912): consider removing this function thanks to rewrite

mlir/include/mlir/Dialect/SPIRV/SPIRVLowering.h

 
 namespace mlir {
 
 /// Type conversion from standard types to SPIR-V types for shader interface.
 ///
 /// For composite types, this converter additionally performs type wrapping to
 /// satisfy shader interface requirements: shader interface types must be
 /// pointers to structs.
-class SPIRVTypeConverter final : public TypeConverter {
+class SPIRVTypeConverter : public TypeConverter {
 public:
   using TypeConverter::TypeConverter;
 
   /// Converts the given standard `type` to SPIR-V correspondence.
   Type convertType(Type type) override;
 
   /// Gets the SPIR-V correspondence for the standard index type.
   static Type getIndexType(MLIRContext *context);

 /// interface/ABI; they convert function parameters to be of SPIR-V allowed
 /// types.
 void populateBuiltinFuncToSPIRVPatterns(MLIRContext *context,
                                         SPIRVTypeConverter &typeConverter,
                                         OwningRewritePatternList &patterns);
 
 namespace spirv {
 class AccessChainOp;
+class FuncOp;
 
 class SPIRVConversionTarget : public ConversionTarget {
 public:
   /// Creates a SPIR-V conversion target for the given target environment.
   static std::unique_ptr<SPIRVConversionTarget> get(TargetEnvAttr targetEnv,
                                                     MLIRContext *context);
 
 private:

 // with AffineMap that has static strides. Extend to handle dynamic strides.
 spirv::AccessChainOp getElementPtr(SPIRVTypeConverter &typeConverter,
                                    MemRefType baseType, Value basePtr,
                                    ArrayRef<Value> indices, Location loc,
                                    OpBuilder &builder);
 
 /// Sets the InterfaceVarABIAttr and EntryPointABIAttr for a function and its
 /// arguments.
-LogicalResult setABIAttrs(FuncOp funcOp, EntryPointABIAttr entryPointInfo,
+LogicalResult setABIAttrs(spirv::FuncOp funcOp,
+                          EntryPointABIAttr entryPointInfo,
                           ArrayRef<InterfaceVarABIAttr> argABIInfo);
 } // namespace spirv
 } // namespace mlir
 
 #endif // MLIR_DIALECT_SPIRV_SPIRVLOWERING_H

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h

 // Get the function that can be used to symbolize an enum value.
 // template <typename EnumClass>
 // Optional<EnumClass> (*)(StringRef) symbolizeEnum();
 #include "mlir/Dialect/SPIRV/SPIRVOpUtils.inc"
 
 } // end namespace spirv
 } // end namespace mlir
 
+namespace llvm {
+
+/// spirv::Function ops hash just like pointers.
+template <>
+struct DenseMapInfo<mlir::spirv::FuncOp> {
+  static mlir::spirv::FuncOp getEmptyKey() {
+    auto pointer = llvm::DenseMapInfo<void *>::getEmptyKey();
+    return mlir::spirv::FuncOp::getFromOpaquePointer(pointer);
+  }
+  static mlir::spirv::FuncOp getTombstoneKey() {
+    auto pointer = llvm::DenseMapInfo<void *>::getTombstoneKey();
+    return mlir::spirv::FuncOp::getFromOpaquePointer(pointer);
+  }
+  static unsigned getHashValue(mlir::spirv::FuncOp val) {
+    return hash_value(val.getAsOpaquePointer());
+  }
+  static bool isEqual(mlir::spirv::FuncOp LHS, mlir::spirv::FuncOp RHS) {
+    return LHS == RHS;
+  }
+};
+
+/// Allow stealing the low bits of spirv::Function ops.
+template <>
+struct PointerLikeTypeTraits<mlir::spirv::FuncOp> {
+public:
+  static inline void *getAsVoidPointer(mlir::spirv::FuncOp I) {
+    return const_cast<void *>(I.getAsOpaquePointer());
+  }
+  static inline mlir::spirv::FuncOp getFromVoidPointer(void *P) {
+    return mlir::spirv::FuncOp::getFromOpaquePointer(P);
+  }
+  static constexpr int NumLowBitsAvailable = 3;
+};
+
+} // namespace llvm
+
 #endif // MLIR_DIALECT_SPIRV_SPIRVOPS_H_

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.td

def SPV_ExecutionModeOp : SPV_Op<"ExecutionMode", [InModuleScope]> {
 
   let results = (outs);
 
   let verifier = [{ return success(); }];
 
   let autogenSerialization = 0;
 
   let builders = [OpBuilder<[{Builder *builder, OperationState &state,
-                              FuncOp function,
+                              spirv::FuncOp function,
                               spirv::ExecutionMode executionMode,
                               ArrayRef<int32_t> params}]>];
 }
 
 // -----
 
 def SPV_LoadOp : SPV_Op<"Load", []> {
   let summary = "Load through a pointer.";

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

 // and module-level operations. The representational form of these ops deviate
 // from the SPIR-V binary format in order to utilize MLIR mechanisms.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef SPIRV_STRUCTURE_OPS
 #define SPIRV_STRUCTURE_OPS
 
+include "mlir/Analysis/CallInterfaces.td"
 include "mlir/Dialect/SPIRV/SPIRVBase.td"
 
+// -----
+
 def SPV_AddressOfOp : SPV_Op<"_address_of", [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
def SPV_AddressOfOp : SPV_Op<"_address_of", [InFunctionScope, NoSideEffect]> {
   let autogenSerialization = 0;
 
   let builders = [OpBuilder<[{Builder *builder, OperationState &state,
                               spirv::GlobalVariableOp var}]>];
 
   let assemblyFormat = "$variable attr-dict `:` type($pointer)";
 }
 
+// -----
+
 def SPV_ConstantOp : SPV_Op<"constant", [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
static spirv::ConstantOp getOne(Type type, Location loc,
                                     OpBuilder *builder);
   }];
 
   let hasOpcode = 0;
 
   let autogenSerialization = 0;
 }
 
+// -----
+
 def SPV_EntryPointOp : SPV_Op<"EntryPoint", [InModuleScope]> {
   let summary = [{
     Declare an entry point, its execution model, and its interface.
   }];
 
   let description = [{
     Execution Model is the execution model for the entry point and its
     static call tree. See Execution Model.
def SPV_EntryPointOp : SPV_Op<"EntryPoint", [InModuleScope]> {
   );
 
   let results = (outs);
 
   let autogenSerialization = 0;
 
   let builders = [OpBuilder<[{Builder *builder, OperationState &state,
                               spirv::ExecutionModel executionModel,
-                              FuncOp function,
+                              spirv::FuncOp function,
                               ArrayRef<Attribute> interfaceVars}]>];
 }
 
+// -----
+
+def SPV_FuncOp : SPV_Op<"func", [
+    DeclareOpInterfaceMethods<CallableOpInterface>,
+    FunctionLike, InModuleScope, IsolatedFromAbove, Symbol
+  ]> {
+  let summary = "Declare or define a function";
+
+  let description = [{
+    This op declares or defines a SPIR-V function using one region, which
+    contains one or more blocks.
+
+    Different from the SPIR-V binary format, this op is not allowed to
+    implicitly capture global values, and all external references must use
+    function arguments or symbol references. This op itself defines a symbol
+    that is unique in the enclosing module op.
+
+    This op itself takes no operands and generates no results. Its region
+    can take zero or more arguments and return zero or one values.
+
+    ### Custom assembly form
+
+    ```
+    spv-function-control ::= "None" | "Inline" | "DontInline" | ...
+    spv-function-op ::= `spv.func` function-signature
+                         spv-function-control region
+    ```
+
+    For example:
+
+    ```
+    spv.func @foo() -> () "None" { ... }
+    spv.func @bar() -> () "Inline|Pure" { ... }
+    ```
+  }];
+
+  let arguments = (ins
+    TypeAttr:$type,
+    StrAttr:$sym_name,
+    SPV_FunctionControlAttr:$function_control
+  );
+
+  let results = (outs);
+
+  let regions = (region AnyRegion:$body);
+
+  let verifier = [{ return success(); }];
+
+  let builders = [OpBuilder<[{
+    Builder *, OperationState &state,
+    StringRef name, FunctionType type,
+    spirv::FunctionControl control = spirv::FunctionControl::None,
+    ArrayRef<NamedAttribute> attrs = {}
+  }]>];
+
+  let hasOpcode = 0;
+
+  let autogenSerialization = 0;
+
+  let extraClassDeclaration = [{
+  private:
+    // This trait needs access to the hooks defined below.
+    friend class OpTrait::FunctionLike<FuncOp>;
+
+    /// Returns the number of arguments. Hook for OpTrait::FunctionLike.
+    unsigned getNumFuncArguments() { return getType().getNumInputs(); }
+
+    /// Returns the number of results. Hook for OpTrait::FunctionLike.
+    unsigned getNumFuncResults() { return getType().getNumResults(); }
+
+    /// Hook for OpTrait::FunctionLike, called after verifying that the 'type'
+    /// attribute is present and checks if it holds a function type. Ensures
+    /// getType, getNumFuncArguments, and getNumFuncResults can be called safely
+    LogicalResult verifyType();
+
+    /// Hook for OpTrait::FunctionLike, called after verifying the function
+    /// type and the presence of the (potentially empty) function body.
+    /// Ensures SPIR-V specific semantics.
+    LogicalResult verifyBody();
+  }];
+}
+
+// -----
 
 def SPV_GlobalVariableOp : SPV_Op<"globalVariable", [InModuleScope, Symbol]> {
   let summary = [{
     Allocate an object in memory at module scope. The object is
     referenced using a symbol name.
   }];
 
   let description = [{
     The variable type must be an OpTypePointer. Its type operand is the type of
     object in memory.
 
     Storage Class is the Storage Class of the memory holding the object. It
     cannot be Generic. It must be the same as the Storage Class operand of
     the variable types. Only those storage classes that are valid at module
     scope (like Input, Output, StorageBuffer, etc.) are valid.
 
     Initializer is optional.  If Initializer is present, it will be
def SPV_GlobalVariableOp : SPV_Op<"globalVariable", [InModuleScope, Symbol]> {
   }];
 
   let arguments = (ins
     TypeAttr:$type,
     StrAttr:$sym_name,
     OptionalAttr<FlatSymbolRefAttr>:$initializer
   );
 
+  let results = (outs);
+
   let builders = [
     OpBuilder<"Builder *builder, OperationState &state, "
       "TypeAttr type, ArrayRef<NamedAttribute> namedAttrs", [{
       state.addAttribute("type", type);
       state.addAttributes(namedAttrs);
     }]>,
     OpBuilder<[{Builder *builder, OperationState &state,
                 Type type, StringRef name, unsigned descriptorSet,
                 unsigned binding}]>,
     OpBuilder<[{Builder *builder, OperationState &state,
                 Type type, StringRef name, spirv::BuiltIn builtin}]>
   ];
 
-  let results = (outs);
-
   let hasOpcode = 0;
 
   let autogenSerialization = 0;
 
   let extraClassDeclaration = [{
     ::mlir::spirv::StorageClass storageClass() {
       return this->type().cast<::mlir::spirv::PointerType>().getStorageClass();
     }
   }];
 }
 
+// -----
+
 def SPV_ModuleOp : SPV_Op<"module",
                           [IsolatedFromAbove,
                            SingleBlockImplicitTerminator<"ModuleEndOp">,
-                           NativeOpTrait<"SymbolTable">]> {
+                           SymbolTable]> {
   let summary = "The top-level op that defines a SPIR-V module";
 
   let description = [{
     This op defines a SPIR-V module using a MLIR region. The region contains
     one block. Module-level operations, including functions definitions,
     are all placed in this block.
 
     Using an op with a region to define a SPIR-V module enables "embedding"
def SPV_ModuleOp : SPV_Op<"module",
 
   let extraClassDeclaration = [{
     Block& getBlock() {
       return this->getOperation()->getRegion(0).front();
     }
   }];
 }
 
+// -----
+
 def SPV_ModuleEndOp : SPV_Op<"_module_end", [InModuleScope, Terminator]> {
   let summary = "The pseudo op that ends a SPIR-V module";
 
   let description = [{
     This op terminates the only block inside a `spv.module`'s only region.
     This op does not have a corresponding SPIR-V instruction and thus will
     not be serialized into the binary format; it is used solely to satisfy
     the structual requirement that an block must be ended with a terminator.
   }];
 
   let arguments = (ins);
 
   let results = (outs);
 
   let assemblyFormat = "attr-dict";
 
   let verifier = [{ return success(); }];
 
   let hasOpcode = 0;
 
   let autogenSerialization = 0;
 }
 
+// -----
+
 def SPV_ReferenceOfOp : SPV_Op<"_reference_of", [NoSideEffect]> {
   let summary = "Reference a specialization constant.";
 
   let description = [{
     Specialization constant 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
def SPV_ReferenceOfOp : SPV_Op<"_reference_of", [NoSideEffect]> {
 
   let hasOpcode = 0;
 
   let autogenSerialization = 0;
 
   let assemblyFormat = "$spec_const attr-dict `:` type($reference)";
 }
 
+// -----
+
 def SPV_SpecConstantOp : SPV_Op<"specConstant", [InModuleScope, Symbol]> {
   let summary = "The op that declares a SPIR-V specialization constant";
 
   let description = [{
     This op declares a SPIR-V scalar specialization constant. SPIR-V has
     multiple constant instructions covering different scalar types:
 
     * `OpSpecConstantTrue` and `OpSpecConstantFalse` for boolean constants
def SPV_SpecConstantOp : SPV_Op<"specConstant", [InModuleScope, Symbol]> {
 
   let results = (outs);
 
   let hasOpcode = 0;
 
   let autogenSerialization = 0;
 }
 
+// -----
+
 #endif // SPIRV_STRUCTURE_OPS

mlir/include/mlir/IR/Function.h

public:
   LogicalResult verify();
 
   /// Erase a single argument at `argIndex`.
   void eraseArgument(unsigned argIndex) { eraseArguments({argIndex}); }
   /// Erases the arguments listed in `argIndices`.
   /// `argIndices` is allowed to have duplicates and can be in any order.
   void eraseArguments(ArrayRef<unsigned> argIndices);
 
-  /// Returns the type of this function.
-  FunctionType getType() {
-    return getAttrOfType<TypeAttr>(getTypeAttrName())
-        .getValue()
-        .cast<FunctionType>();
-  }
-
-  /// Change the type of this function in place. This is an extremely dangerous
-  /// operation and it is up to the caller to ensure that this is legal for this
-  /// function, and to restore invariants:
-  ///  - the entry block args must be updated to match the function params.
-  ///  - the argument/result attributes may need an update: if the new type has
-  ///  less parameters we drop the extra attributes, if there are more
-  ///  parameters they won't have any attributes.
-  void setType(FunctionType newType) {
-    SmallVector<char, 16> nameBuf;
-    auto oldType = getType();
-    for (int i = newType.getNumInputs(), e = oldType.getNumInputs(); i < e;
-         i++) {
-      removeAttr(getArgAttrName(i, nameBuf));
-    }
-    for (int i = newType.getNumResults(), e = oldType.getNumResults(); i < e;
-         i++) {
-      removeAttr(getResultAttrName(i, nameBuf));
-    }
-    setAttr(getTypeAttrName(), TypeAttr::get(newType));
-  }
-
   /// Create a deep copy of this function and all of its blocks, remapping
   /// any operands that use values outside of the function using the map that is
   /// provided (leaving them alone if no entry is present). If the mapper
   /// contains entries for function arguments, these arguments are not included
   /// in the new function. Replaces references to cloned sub-values with the
   /// corresponding value that is copied, and adds those mappings to the mapper.
   FuncOp clone(BlockAndValueMapping &mapper);
   FuncOp clone();
 
   /// Clone the internal blocks and attributes from this function into dest. Any
   /// cloned blocks are appended to the back of dest. This function asserts that
   /// the attributes of the current function and dest are compatible.
   void cloneInto(FuncOp dest, BlockAndValueMapping &mapper);
 
   //===--------------------------------------------------------------------===//
-  // Body Handling
-  //===--------------------------------------------------------------------===//
-
-  /// Add an entry block to an empty function, and set up the block arguments
-  /// to match the signature of the function. The newly inserted entry block is
-  /// returned.
-  Block *addEntryBlock();
-
-  /// Add a normal block to the end of the function's block list. The function
-  /// should at least already have an entry block.
-  Block *addBlock();
-
-  //===--------------------------------------------------------------------===//
   // CallableOpInterface
   //===--------------------------------------------------------------------===//
 
   /// Returns the region on the current operation that is callable. This may
   /// return null in the case of an external callable object, e.g. an external
   /// function.
   Region *getCallableRegion() { return isExternal() ? nullptr : &getBody(); }
 

mlir/include/mlir/IR/FunctionSupport.h

public:
 
   bool empty() { return getBody().empty(); }
   void push_back(Block *block) { getBody().push_back(block); }
   void push_front(Block *block) { getBody().push_front(block); }
 
   Block &back() { return getBody().back(); }
   Block &front() { return getBody().front(); }
 
+  /// Add an entry block to an empty function, and set up the block arguments
+  /// to match the signature of the function. The newly inserted entry block
+  /// is returned.
+  ///
+  /// Note that the concrete class must define a method with the same name to
+  /// hide this one if the concrete class does not use FunctionType for the
+  /// function type under the hood.
+  Block *addEntryBlock();
+
+  /// Add a normal block to the end of the function's block list. The function
+  /// should at least already have an entry block.
+  Block *addBlock();
+
   /// Hook for concrete ops to verify the contents of the body. Called as a
   /// part of trait verification, after type verification and ensuring that a
   /// region exists.
   LogicalResult verifyBody();
 
   //===--------------------------------------------------------------------===//
   // Type Attribute Handling
   //===--------------------------------------------------------------------===//
 
   /// Return the name of the attribute used for function types.
   static StringRef getTypeAttrName() { return ::mlir::impl::getTypeAttrName(); }
 
   TypeAttr getTypeAttr() {
     return this->getOperation()->template getAttrOfType<TypeAttr>(
         getTypeAttrName());
   }
 
+  /// Return the type of this function.
+  ///
+  /// Note that the concrete class must define a method with the same name to
+  /// hide this one if the concrete class does not use FunctionType for the
+  /// function type under the hood.
+  FunctionType getType() {
+    return getTypeAttr().getValue().template cast<FunctionType>();
+  }
+
   bool isTypeAttrValid() {
     auto typeAttr = getTypeAttr();
     if (!typeAttr)
       return false;
     return typeAttr.getValue() != Type{};
   }
 
+  /// Change the type of this function in place. This is an extremely dangerous
+  /// operation and it is up to the caller to ensure that this is legal for this
+  /// function, and to restore invariants:
+  ///  - the entry block args must be updated to match the function params.
+  ///  - the argument/result attributes may need an update: if the new type
+  ///    has less parameters we drop the extra attributes, if there are more
+  ///    parameters they won't have any attributes.
+  ///
+  /// Note that the concrete class must define a method with the same name to
+  /// hide this one if the concrete class does not use FunctionType for the
+  /// function type under the hood.
+  void setType(FunctionType newType);
+
   //===--------------------------------------------------------------------===//
   // Argument Handling
   //===--------------------------------------------------------------------===//
 
   unsigned getNumArguments() {
     return static_cast<ConcreteType *>(this)->getNumFuncArguments();
   }
 
LogicalResult FunctionLike<ConcreteType>::verifyTrait(Operation *op) {
   // Check that the op has exactly one region for the body.
   if (op->getNumRegions() != 1)
     return funcOp.emitOpError("expects one region");
 
   return funcOp.verifyBody();
 }
 
 //===----------------------------------------------------------------------===//
+// Function Body.
+//===----------------------------------------------------------------------===//
+
+template <typename ConcreteType>
+Block *FunctionLike<ConcreteType>::addEntryBlock() {
+  assert(empty() && "function already has an entry block");
+  auto *entry = new Block();
+  push_back(entry);
+  entry->addArguments(getType().getInputs());
+  return entry;
+}
+
+template <typename ConcreteType>
+Block *FunctionLike<ConcreteType>::addBlock() {
+  assert(!empty() && "function should at least have an entry block");
+  push_back(new Block());
+  return &back();
+}
+
+//===----------------------------------------------------------------------===//
+// Function Type Attribute.
+//===----------------------------------------------------------------------===//
+
+template <typename ConcreteType>
+void FunctionLike<ConcreteType>::setType(FunctionType newType) {
+  SmallVector<char, 16> nameBuf;
+  auto oldType = getType();
+  auto *concreteOp = static_cast<ConcreteType *>(this);
+
+  for (int i = newType.getNumInputs(), e = oldType.getNumInputs(); i < e; i++)
+    concreteOp->removeAttr(getArgAttrName(i, nameBuf));
+  for (int i = newType.getNumResults(), e = oldType.getNumResults(); i < e; i++)
+    concreteOp->removeAttr(getResultAttrName(i, nameBuf));
+  concreteOp->setAttr(getTypeAttrName(), TypeAttr::get(newType));
+}
+
+//===----------------------------------------------------------------------===//
 // Function Argument Attribute.
 //===----------------------------------------------------------------------===//
 
 /// Set the attributes held by the argument at 'index'.
 template <typename ConcreteType>
 void FunctionLike<ConcreteType>::setArgAttrs(
     unsigned index, ArrayRef<NamedAttribute> attributes) {
   assert(index < getNumArguments() && "invalid argument number");

mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp

PatternMatchResult WorkGroupSizeConversion::matchAndRewrite(
   return matchSuccess();
 }
 
 //===----------------------------------------------------------------------===//
 // GPUFuncOp
 //===----------------------------------------------------------------------===//
 
 // Legalizes a GPU function as an entry SPIR-V function.
-static FuncOp
+static spirv::FuncOp
 lowerAsEntryFunction(gpu::GPUFuncOp funcOp, SPIRVTypeConverter &typeConverter,
                      ConversionPatternRewriter &rewriter,
                      spirv::EntryPointABIAttr entryPointInfo,
                      ArrayRef<spirv::InterfaceVarABIAttr> argABIInfo) {
   auto fnType = funcOp.getType();
   if (fnType.getNumResults()) {
     funcOp.emitError("SPIR-V lowering only supports entry functions"
                      "with no return values right now");
lowerAsEntryFunction(gpu::GPUFuncOp funcOp, SPIRVTypeConverter &typeConverter,
   // LowerABIAttributesPass.
   TypeConverter::SignatureConversion signatureConverter(fnType.getNumInputs());
   {
     for (auto argType : enumerate(funcOp.getType().getInputs())) {
       auto convertedType = typeConverter.convertType(argType.value());
       signatureConverter.addInputs(argType.index(), convertedType);
     }
   }
-  auto newFuncOp = rewriter.create<FuncOp>(
+  auto newFuncOp = rewriter.create<spirv::FuncOp>(
       funcOp.getLoc(), funcOp.getName(),
       rewriter.getFunctionType(signatureConverter.getConvertedTypes(),
-                               llvm::None),
-      ArrayRef<NamedAttribute>());
+                               llvm::None));
   for (const auto &namedAttr : funcOp.getAttrs()) {
     if (namedAttr.first.is(impl::getTypeAttrName()) ||
         namedAttr.first.is(SymbolTable::getSymbolAttrName()))
       continue;
     newFuncOp.setAttr(namedAttr.first, namedAttr.second);
   }
   rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(),
                               newFuncOp.end());
argABI.push_back(spirv::getInterfaceVarABIAttr(
         0, argNum, spirv::StorageClass::StorageBuffer, rewriter.getContext()));
   }
 
   auto entryPointAttr = spirv::lookupEntryPointABI(funcOp);
   if (!entryPointAttr) {
     funcOp.emitRemark("match failure: missing 'spv.entry_point_abi' attribute");
     return matchFailure();
   }
-  FuncOp newFuncOp = lowerAsEntryFunction(funcOp, typeConverter, rewriter,
-                                          entryPointAttr, argABI);
+  spirv::FuncOp newFuncOp = lowerAsEntryFunction(
+      funcOp, typeConverter, rewriter, entryPointAttr, argABI);
   if (!newFuncOp)
     return matchFailure();
   newFuncOp.removeAttr(Identifier::get(gpu::GPUDialect::getKernelFuncAttrName(),
                                        rewriter.getContext()));
   return matchSuccess();
 }
 
 //===----------------------------------------------------------------------===//

mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRVPass.cpp

void GPUToSPIRVPass::runOnModule() {
 
   SPIRVTypeConverter typeConverter;
   OwningRewritePatternList patterns;
   populateGPUToSPIRVPatterns(context, typeConverter, patterns);
   populateStandardToSPIRVPatterns(context, typeConverter, patterns);
 
   std::unique_ptr<ConversionTarget> target = spirv::SPIRVConversionTarget::get(
       spirv::lookupTargetEnvOrDefault(module), context);
-  target->addDynamicallyLegalOp<FuncOp>(
-      [&](FuncOp op) { return typeConverter.isSignatureLegal(op.getType()); });
 
   if (failed(applyFullConversion(kernelModules, *target, patterns,
                                  &typeConverter))) {
     return signalPassFailure();
   }
 }
 
 std::unique_ptr<OpPassBase<ModuleOp>> mlir::createConvertGPUToSPIRVPass() {
   return std::make_unique<GPUToSPIRVPass>();
 }
 
 static PassRegistration<GPUToSPIRVPass>
     pass("convert-gpu-to-spirv", "Convert GPU dialect to SPIR-V dialect");

mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRVPass.cpp

void ConvertStandardToSPIRVPass::runOnModule() {
   ModuleOp module = getModule();
 
   SPIRVTypeConverter typeConverter;
   OwningRewritePatternList patterns;
   populateStandardToSPIRVPatterns(context, typeConverter, patterns);
   populateBuiltinFuncToSPIRVPatterns(context, typeConverter, patterns);
 
   std::unique_ptr<ConversionTarget> target = spirv::SPIRVConversionTarget::get(
       spirv::lookupTargetEnvOrDefault(module), context);
-  target->addDynamicallyLegalOp<FuncOp>(
-      [&](FuncOp op) { return typeConverter.isSignatureLegal(op.getType()); });
 
   if (failed(applyPartialConversion(module, *target, patterns))) {
     return signalPassFailure();
   }
 }
 
 std::unique_ptr<OpPassBase<ModuleOp>> mlir::createConvertStandardToSPIRVPass() {
   return std::make_unique<ConvertStandardToSPIRVPass>();
 }
 
 static PassRegistration<ConvertStandardToSPIRVPass>
     pass("convert-std-to-spirv", "Convert Standard Ops to SPIR-V dialect");

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

 /// This class defines the interface for inlining within the SPIR-V dialect.
 struct SPIRVInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   /// Returns true if the given region 'src' can be inlined into the region
   /// 'dest' that is attached to an operation registered to the current dialect.
   bool isLegalToInline(Region *dest, Region *src,
                        BlockAndValueMapping &) const final {
-    // Return true here when inlining into spv.selection and spv.loop
-    // operations.
+    // Return true here when inlining into spv.func, spv.selection, and
+    // spv.loop operations.
     auto op = dest->getParentOp();
-    return isa<spirv::SelectionOp>(op) || isa<spirv::LoopOp>(op);
+    return isa<spirv::FuncOp>(op) || isa<spirv::SelectionOp>(op) ||
+           isa<spirv::LoopOp>(op);
   }
 
   /// Returns true if the given operation 'op', that is registered to this
   /// dialect, can be inlined into the region 'dest' that is attached to an
   /// operation registered to the current dialect.
   bool isLegalToInline(Operation *op, Region *dest,
                        BlockAndValueMapping &) const final {
     // TODO(antiagainst): Enable inlining structured control flows with return.

mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp

matchAndRewrite(FuncOp funcOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 } // namespace
 
 PatternMatchResult
 FuncOpConversion::matchAndRewrite(FuncOp funcOp, ArrayRef<Value> operands,
                                   ConversionPatternRewriter &rewriter) const {
   auto fnType = funcOp.getType();
+  // TODO(antiagainst): support converting functions with one result.
   if (fnType.getNumResults())
     return matchFailure();
 
   TypeConverter::SignatureConversion signatureConverter(fnType.getNumInputs());
   for (auto argType : enumerate(funcOp.getType().getInputs())) {
     auto convertedType = typeConverter.convertType(argType.value());
     if (!convertedType)
       return matchFailure();
     signatureConverter.addInputs(argType.index(), convertedType);
   }
 
-  rewriter.updateRootInPlace(funcOp, [&] {
-    funcOp.setType(rewriter.getFunctionType(
-        signatureConverter.getConvertedTypes(), llvm::None));
-    rewriter.applySignatureConversion(&funcOp.getBody(), signatureConverter);
-  });
-
+  // Create the converted spv.func op.
+  auto newFuncOp = rewriter.create<spirv::FuncOp>(
+      funcOp.getLoc(), funcOp.getName(),
+      rewriter.getFunctionType(signatureConverter.getConvertedTypes(),
+                               llvm::None));
+
+  // Copy over all attributes other than the function name and type.
+  for (const auto &namedAttr : funcOp.getAttrs()) {
+    if (!namedAttr.first.is(impl::getTypeAttrName()) &&
+        !namedAttr.first.is(SymbolTable::getSymbolAttrName()))
+      newFuncOp.setAttr(namedAttr.first, namedAttr.second);
+  }
+
+  rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(),
+                              newFuncOp.end());
+  rewriter.applySignatureConversion(&newFuncOp.getBody(), signatureConverter);
+  rewriter.eraseOp(funcOp);
   return matchSuccess();
 }
 
 void mlir::populateBuiltinFuncToSPIRVPatterns(
     MLIRContext *context, SPIRVTypeConverter &typeConverter,
     OwningRewritePatternList &patterns) {
   patterns.insert<FuncOpConversion>(context, typeConverter);
 }
spirv::AccessChainOp mlir::spirv::getElementPtr(
   if (failed(getStridesAndOffset(baseType, strides, offset)) ||
       llvm::is_contained(strides, MemRefType::getDynamicStrideOrOffset())) {
     return nullptr;
   }
 
   auto indexType = typeConverter.getIndexType(builder.getContext());
 
   Value ptrLoc = nullptr;
-  assert(indices.size() == strides.size());
+  assert(indices.size() == strides.size() &&
+         "must provide indices for all dimensions");
   for (auto index : enumerate(indices)) {
     Value strideVal = builder.create<spirv::ConstantOp>(
         loc, indexType, IntegerAttr::get(indexType, strides[index.index()]));
     Value update = builder.create<spirv::IMulOp>(loc, strideVal, index.value());
     ptrLoc =
         (ptrLoc ? builder.create<spirv::IAddOp>(loc, ptrLoc, update).getResult()
                 : update);
   }
   SmallVector<Value, 2> linearizedIndices;
   // Add a '0' at the start to index into the struct.
   linearizedIndices.push_back(builder.create<spirv::ConstantOp>(
       loc, indexType, IntegerAttr::get(indexType, 0)));
   linearizedIndices.push_back(ptrLoc);
   return builder.create<spirv::AccessChainOp>(loc, basePtr, linearizedIndices);
 }
 
 //===----------------------------------------------------------------------===//
 // Set ABI attributes for lowering entry functions.
 //===----------------------------------------------------------------------===//
 
 LogicalResult
-mlir::spirv::setABIAttrs(FuncOp funcOp, spirv::EntryPointABIAttr entryPointInfo,
+mlir::spirv::setABIAttrs(spirv::FuncOp funcOp,
+                         spirv::EntryPointABIAttr entryPointInfo,
                          ArrayRef<spirv::InterfaceVarABIAttr> argABIInfo) {
   // Set the attributes for argument and the function.
   StringRef argABIAttrName = spirv::getInterfaceVarABIAttrName();
   for (auto argIndex : llvm::seq<unsigned>(0, funcOp.getNumArguments())) {
     funcOp.setArgAttr(argIndex, argABIAttrName, argABIInfo[argIndex]);
   }
   funcOp.setAttr(spirv::getEntryPointABIAttrName(), entryPointInfo);
   return success();

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
 
 #include "mlir/Analysis/CallInterfaces.h"
 #include "mlir/Dialect/CommonFolders.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Function.h"
+#include "mlir/IR/FunctionImplementation.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/StandardTypes.h"
 #include "mlir/Support/Functional.h"
 #include "mlir/Support/StringExtras.h"
 #include "llvm/ADT/bit.h"
 

 static bool isDirectInModuleLikeOp(Operation *op) {
   return op && op->hasTrait<OpTrait::SymbolTable>();
 }
 
 static LogicalResult extractValueFromConstOp(Operation *op, int32_t &value) {
   auto constOp = dyn_cast_or_null<spirv::ConstantOp>(op);
   if (!constOp) {
     return failure();
   }
   auto valueAttr = constOp.value();
   auto integerValueAttr = valueAttr.dyn_cast<IntegerAttr>();
   if (!integerValueAttr) {
     return failure();
   }
   value = integerValueAttr.getInt();
   return success();
 }

 }
 
 //===----------------------------------------------------------------------===//
 // spv.EntryPoint
 //===----------------------------------------------------------------------===//
 
 void spirv::EntryPointOp::build(Builder *builder, OperationState &state,
                                 spirv::ExecutionModel executionModel,
-                                FuncOp function,
+                                spirv::FuncOp function,
                                 ArrayRef<Attribute> interfaceVars) {
   build(builder, state,
         builder->getI32IntegerAttr(static_cast<int32_t>(executionModel)),
         builder->getSymbolRefAttr(function),
         builder->getArrayAttr(interfaceVars));
 }
 
 static ParseResult parseEntryPointOp(OpAsmParser &parser,
static LogicalResult verify(spirv::EntryPointOp entryPointOp) {
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // spv.ExecutionMode
 //===----------------------------------------------------------------------===//
 
 void spirv::ExecutionModeOp::build(Builder *builder, OperationState &state,
-                                   FuncOp function,
+                                   spirv::FuncOp function,
                                    spirv::ExecutionMode executionMode,
                                    ArrayRef<int32_t> params) {
   build(builder, state, builder->getSymbolRefAttr(function),
         builder->getI32IntegerAttr(static_cast<int32_t>(executionMode)),
         builder->getI32ArrayAttr(params));
 }
 
 static ParseResult parseExecutionModeOp(OpAsmParser &parser,
while (!parser.parseOptionalComma()) {
     values.push_back(value.cast<IntegerAttr>().getInt());
   }
   state.addAttribute(kValuesAttrName,
                      parser.getBuilder().getI32ArrayAttr(values));
   return success();
 }
 
 static void print(spirv::ExecutionModeOp execModeOp, OpAsmPrinter &printer) {
-  printer << spirv::ExecutionModeOp::getOperationName() << " @"
-          << execModeOp.fn() << " \""
-          << stringifyExecutionMode(execModeOp.execution_mode()) << "\"";
+  printer << spirv::ExecutionModeOp::getOperationName() << " ";
+  printer.printSymbolName(execModeOp.fn());
+  printer << " \"" << stringifyExecutionMode(execModeOp.execution_mode())
+          << "\"";
   auto values = execModeOp.values();
   if (!values.size())
     return;
   printer << ", ";
   interleaveComma(values, printer, [&](Attribute a) {
     printer << a.cast<IntegerAttr>().getInt();
   });
 }
 
 //===----------------------------------------------------------------------===//
+// spv.func
+//===----------------------------------------------------------------------===//
+
+static ParseResult parseFuncOp(OpAsmParser &parser, OperationState &state) {
+  SmallVector<OpAsmParser::OperandType, 4> entryArgs;
+  SmallVector<SmallVector<NamedAttribute, 2>, 4> argAttrs;
+  SmallVector<SmallVector<NamedAttribute, 2>, 4> resultAttrs;
+  SmallVector<Type, 4> argTypes;
+  SmallVector<Type, 4> resultTypes;
+  auto &builder = parser.getBuilder();
+
+  // Parse the name as a symbol.
+  StringAttr nameAttr;
+  if (parser.parseSymbolName(nameAttr, SymbolTable::getSymbolAttrName(),
+                             state.attributes))
+    return failure();
+
+  // Parse the function signature.
+  bool isVariadic = false;
+  if (impl::parseFunctionSignature(parser, /*allowVariadic=*/false, entryArgs,
+                                   argTypes, argAttrs, isVariadic, resultTypes,
+                                   resultAttrs))
+    return failure();
+
+  auto fnType = builder.getFunctionType(argTypes, resultTypes);
+  state.addAttribute(impl::getTypeAttrName(), TypeAttr::get(fnType));
+
+  // Parse the optional function control keyword.
+  spirv::FunctionControl fnControl;
+  if (parseEnumAttribute(fnControl, parser, state))
+    return failure();
+
+  // If additional attributes are present, parse them.
+  if (parser.parseOptionalAttrDictWithKeyword(state.attributes))
+    return failure();
+
+  // Add the attributes to the function arguments.
+  assert(argAttrs.size() == argTypes.size());
+  assert(resultAttrs.size() == resultTypes.size());
+  impl::addArgAndResultAttrs(builder, state, argAttrs, resultAttrs);
+
+  // Parse the optional function body.
+  auto *body = state.addRegion();
+  return parser.parseOptionalRegion(
+      *body, entryArgs, entryArgs.empty() ? ArrayRef<Type>() : argTypes);
+}
+
+static void print(spirv::FuncOp fnOp, OpAsmPrinter &printer) {
+  // Print function name, signature, and control.
+  printer << spirv::FuncOp::getOperationName() << " ";
+  printer.printSymbolName(fnOp.sym_name());
+  auto fnType = fnOp.getType();
+  impl::printFunctionSignature(printer, fnOp, fnType.getInputs(),
+                               /*isVariadic=*/false, fnType.getResults());
+  printer << " \"" << spirv::stringifyFunctionControl(fnOp.function_control())
+          << "\"";
+  impl::printFunctionAttributes(
+      printer, fnOp, fnType.getNumInputs(), fnType.getNumResults(),
+      {spirv::attributeName<spirv::FunctionControl>()});
+
+  // Print the body if this is not an external function.
+  Region &body = fnOp.body();
+  if (!body.empty())
+    printer.printRegion(body, /*printEntryBlockArgs=*/false,
+                        /*printBlockTerminators=*/true);
+}
+
+LogicalResult spirv::FuncOp::verifyType() {
+  auto type = getTypeAttr().getValue();
+  if (!type.isa<FunctionType>())
+    return emitOpError("requires '" + getTypeAttrName() +
+                       "' attribute of function type");
+  if (getType().getNumResults() > 1)
+    return emitOpError("cannot have more than one result");
+  return success();
+}
+
+LogicalResult spirv::FuncOp::verifyBody() {
+  FunctionType fnType = getType();
+
+  auto walkResult = walk([fnType](Operation *op) -> WalkResult {
+    if (auto retOp = dyn_cast<spirv::ReturnOp>(op)) {
+      if (fnType.getNumResults() != 0)
+        return retOp.emitOpError("cannot be used in functions returning value");
+    } else if (auto retOp = dyn_cast<spirv::ReturnValueOp>(op)) {
+      if (fnType.getNumResults() != 1)
+        return retOp.emitOpError(
+                   "returns 1 value but enclosing function requires ")
+               << fnType.getNumResults() << " results";
+
+      auto retOperandType = retOp.value().getType();
+      auto fnResultType = fnType.getResult(0);
+      if (retOperandType != fnResultType)
+        return retOp.emitOpError(" return value's type (")
+               << retOperandType << ") mismatch with function's result type ("
+               << fnResultType << ")";
+    }
+    return WalkResult::advance();
+  });
+
+  // TODO(antiagainst): verify other bits like linkage type.
+
+  return failure(walkResult.wasInterrupted());
+}
+
+void spirv::FuncOp::build(Builder *builder, OperationState &state,
+                          StringRef name, FunctionType type,
+                          spirv::FunctionControl control,
+                          ArrayRef<NamedAttribute> attrs) {
+  state.addAttribute(SymbolTable::getSymbolAttrName(),
+                     builder->getStringAttr(name));
+  state.addAttribute(getTypeAttrName(), TypeAttr::get(type));
+  state.addAttribute(
+      spirv::attributeName<spirv::FunctionControl>(),
+      builder->getI32IntegerAttr(static_cast<uint32_t>(control)));
+  state.attributes.append(attrs.begin(), attrs.end());
+  state.addRegion();
+}
+
+// CallableOpInterface
+Region *spirv::FuncOp::getCallableRegion() {
+  return isExternal() ? nullptr : &body();
+}
+
+// CallableOpInterface
+ArrayRef<Type> spirv::FuncOp::getCallableResults() {
+  return getType().getResults();
+}
+
+//===----------------------------------------------------------------------===//
 // spv.FunctionCall
 //===----------------------------------------------------------------------===//
 
 static LogicalResult verify(spirv::FunctionCallOp functionCallOp) {
   auto fnName = functionCallOp.callee();
 
-  auto funcOp = dyn_cast_or_null<FuncOp>(SymbolTable::lookupNearestSymbolFrom(
+  auto funcOp =
+      dyn_cast_or_null<spirv::FuncOp>(SymbolTable::lookupNearestSymbolFrom(
           functionCallOp.getParentOp(), fnName));
   if (!funcOp) {
     return functionCallOp.emitOpError("callee function '")
            << fnName << "' not found in nearest symbol table";
   }
 
   auto functionType = funcOp.getType();
 
   if (functionCallOp.getNumResults() > 1) {
printer.printRegion(moduleOp.body(), /*printEntryBlockArgs=*/false,
                       /*printBlockTerminators=*/false);
   printer.printOptionalAttrDictWithKeyword(moduleOp.getAttrs(), elidedAttrs);
 }
 
 static LogicalResult verify(spirv::ModuleOp moduleOp) {
   auto &op = *moduleOp.getOperation();
   auto *dialect = op.getDialect();
   auto &body = op.getRegion(0).front();
-  DenseMap<std::pair<FuncOp, spirv::ExecutionModel>, spirv::EntryPointOp>
+  DenseMap<std::pair<spirv::FuncOp, spirv::ExecutionModel>, spirv::EntryPointOp>
       entryPoints;
   SymbolTable table(moduleOp);
 
   for (auto &op : body) {
-    if (op.getDialect() == dialect) {
-      // For EntryPoint op, check that the function and execution model is not
-      // duplicated in EntryPointOps. Also verify that the interface specified
-      // comes from globalVariables here to make this check cheaper.
-      if (auto entryPointOp = dyn_cast<spirv::EntryPointOp>(op)) {
-        auto funcOp = table.lookup<FuncOp>(entryPointOp.fn());
-        if (!funcOp) {
-          return entryPointOp.emitError("function '")
-                 << entryPointOp.fn() << "' not found in 'spv.module'";
-        }
-        if (auto interface = entryPointOp.interface()) {
-          for (Attribute varRef : interface) {
-            auto varSymRef = varRef.dyn_cast<FlatSymbolRefAttr>();
-            if (!varSymRef) {
-              return entryPointOp.emitError(
-                         "expected symbol reference for interface "
-                         "specification instead of '")
-                     << varRef;
-            }
-            auto variableOp =
-                table.lookup<spirv::GlobalVariableOp>(varSymRef.getValue());
-            if (!variableOp) {
-              return entryPointOp.emitError("expected spv.globalVariable "
-                                            "symbol reference instead of'")
-                     << varSymRef << "'";
-            }
-          }
-        }
+    if (op.getDialect() != dialect)
+      return op.emitError("'spv.module' can only contain spv.* ops");
+
+      auto funcOp = table.lookup<spirv::FuncOp>(entryPointOp.fn());
 
-        auto key = std::pair<FuncOp, spirv::ExecutionModel>(
-            funcOp, entryPointOp.execution_model());
-        auto entryPtIt = entryPoints.find(key);
-        if (entryPtIt != entryPoints.end()) {
-          return entryPointOp.emitError("duplicate of a previous EntryPointOp");
-        }
-        entryPoints[key] = entryPointOp;
-      }
-      continue;
-    }
-
-    auto funcOp = dyn_cast<FuncOp>(op);
-    if (!funcOp)
-      return op.emitError("'spv.module' can only contain func and spv.* ops");
-
+      auto key = std::pair<spirv::FuncOp, spirv::ExecutionModel>(
+    } else if (auto funcOp = dyn_cast<spirv::FuncOp>(op)) {
       if (funcOp.isExternal())
         return op.emitError("'spv.module' cannot contain external functions");
 
+      // TODO(antiagainst): move this check to spv.func.
       for (auto &block : funcOp)
         for (auto &op : block) {
-        if (op.getDialect() == dialect)
-          continue;
-
-        if (isa<FuncOp>(op))
-          return op.emitError("'spv.module' cannot contain nested functions");
-
+          if (op.getDialect() != dialect)
             return op.emitError(
                 "functions in 'spv.module' can only contain spv.* ops");
         }
     }
+  }
 
   // Verify capabilities. ODS already guarantees that we have an array of
   // string attributes.
   if (auto caps = moduleOp.getAttrOfType<ArrayAttr>("capabilities")) {
     for (auto cap : caps.getValue()) {
       auto capStr = cap.cast<StringAttr>().getValue();
       if (!spirv::symbolizeCapability(capStr))
         return moduleOp.emitOpError("uses unknown capability: ") << capStr;
static LogicalResult verify(spirv::ReferenceOfOp referenceOfOp) {
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // spv.Return
 //===----------------------------------------------------------------------===//
 
 static LogicalResult verify(spirv::ReturnOp returnOp) {
-  auto funcOp = returnOp.getParentOfType<FuncOp>();
-  auto numOutputs = funcOp.getType().getNumResults();
-  if (numOutputs != 0)
-    return returnOp.emitOpError("cannot be used in functions returning value")
-           << (numOutputs > 1 ? "s" : "");
-
+  // Verification is performed in spv.func op.
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // spv.ReturnValue
 //===----------------------------------------------------------------------===//
 
 static LogicalResult verify(spirv::ReturnValueOp retValOp) {
-  auto funcOp = retValOp.getParentOfType<FuncOp>();
-  auto numFnResults = funcOp.getType().getNumResults();
-  if (numFnResults != 1)
-    return retValOp.emitOpError(
-               "returns 1 value but enclosing function requires ")
-           << numFnResults << " results";
-
-  auto operandType = retValOp.value().getType();
-  auto fnResultType = funcOp.getType().getResult(0);
-  if (operandType != fnResultType)
-    return retValOp.emitOpError(" return value's type (")
-           << operandType << ") mismatch with function's result type ("
-           << fnResultType << ")";
-
+  // Verification is performed in spv.func op.
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // spv.Select
 //===----------------------------------------------------------------------===//
 
 void spirv::SelectOp::build(Builder *builder, OperationState &state, Value cond,
                             Value trueValue, Value falseValue) {
   build(builder, state, trueValue.getType(), cond, trueValue, falseValue);

mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp

 
 /// Extracts the opcode from the given first word of a SPIR-V instruction.
 static inline spirv::Opcode extractOpcode(uint32_t word) {
   return static_cast<spirv::Opcode>(word & 0xffff);
 }
 
 /// Returns true if the given `block` is a function entry block.
 static inline bool isFnEntryBlock(Block *block) {
-  return block->isEntryBlock() && isa_and_nonnull<FuncOp>(block->getParentOp());
+  return block->isEntryBlock() &&
+         isa_and_nonnull<spirv::FuncOp>(block->getParentOp());
 }
 
 namespace {
 /// A struct for containing a header block's merge and continue targets.
 ///
 /// This struct is used to track original structured control flow info from
 /// SPIR-V blob. This info will be used to create spv.selection/spv.loop
 /// later.
private:
   LogicalResult processDecoration(ArrayRef<uint32_t> words);
 
   // Processes an OpMemberDecorate instruction.
   LogicalResult processMemberDecoration(ArrayRef<uint32_t> words);
 
   /// Processes an OpMemberName instruction.
   LogicalResult processMemberName(ArrayRef<uint32_t> words);
 
-  /// Gets the FuncOp associated with a result <id> of OpFunction.
-  FuncOp getFunction(uint32_t id) { return funcMap.lookup(id); }
+  /// Gets the function op associated with a result <id> of OpFunction.
+  spirv::FuncOp getFunction(uint32_t id) { return funcMap.lookup(id); }
 
   /// Processes the SPIR-V function at the current `offset` into `binary`.
   /// The operands to the OpFunction instruction is passed in as ``operands`.
   /// This method processes each instruction inside the function and dispatches
   /// them to their handler method accordingly.
   LogicalResult processFunction(ArrayRef<uint32_t> operands);
 
   /// Processes OpFunctionEnd and finalizes function. This wires up block
private:
 
   // TODO(antiagainst): create Location subclass for binary blob
   Location unknownLoc;
 
   /// The SPIR-V ModuleOp.
   Optional<spirv::ModuleOp> module;
 
   /// The current function under construction.
-  Optional<FuncOp> curFunction;
+  Optional<spirv::FuncOp> curFunction;
 
   /// The current block under construction.
   Block *curBlock = nullptr;
 
   OpBuilder opBuilder;
 
   /// The list of capabilities used by the module.
   llvm::SmallSetVector<spirv::Capability, 4> capabilities;
private:
 
   // Result <id> to variable mapping.
   DenseMap<uint32_t, spirv::SpecConstantOp> specConstMap;
 
   // Result <id> to variable mapping.
   DenseMap<uint32_t, spirv::GlobalVariableOp> globalVariableMap;
 
   // Result <id> to function mapping.
-  DenseMap<uint32_t, FuncOp> funcMap;
+  DenseMap<uint32_t, spirv::FuncOp> funcMap;
 
   // Result <id> to block mapping.
   DenseMap<uint32_t, Block *> blockMap;
 
   // Header block to its merge (and continue) target mapping.
   BlockMergeInfoMap blockMergeInfo;
 
   // Block to its phi (block argument) mapping.
LogicalResult Deserializer::processFunction(ArrayRef<uint32_t> operands) {
   if ((isVoidType(resultType) && functionType.getNumResults() != 0) ||
       (functionType.getNumResults() == 1 &&
        functionType.getResult(0) != resultType)) {
     return emitError(unknownLoc, "mismatch in function type ")
            << functionType << " and return type " << resultType << " specified";
   }
 
   std::string fnName = getFunctionSymbol(operands[1]);
-  auto funcOp = opBuilder.create<FuncOp>(unknownLoc, fnName, functionType,
-                                         ArrayRef<NamedAttribute>());
+  auto funcOp =
+      opBuilder.create<spirv::FuncOp>(unknownLoc, fnName, functionType);
   curFunction = funcMap[operands[1]] = funcOp;
   LLVM_DEBUG(llvm::dbgs() << "-- start function " << fnName << " (type = "
                           << fnType << ", id = " << operands[1] << ") --\n");
   auto *entryBlock = funcOp.addEntryBlock();
   LLVM_DEBUG(llvm::dbgs() << "[block] created entry block " << entryBlock
                           << "\n");
 
   // Parse the op argument instructions

mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp

private:
   /// OpUndef be generated at module level. The serialization generates an
   /// OpUndef for each type needed at module level.
   LogicalResult processUndefOp(spirv::UndefOp op);
 
   /// Emit OpName for the given `resultID`.
   LogicalResult processName(uint32_t resultID, StringRef name);
 
   /// Processes a SPIR-V function op.
-  LogicalResult processFuncOp(FuncOp op);
+  LogicalResult processFuncOp(spirv::FuncOp op);
 
   LogicalResult processVariableOp(spirv::VariableOp op);
 
   /// Process a SPIR-V GlobalVariableOp
   LogicalResult processGlobalVariableOp(spirv::GlobalVariableOp varOp);
 
   /// Process attributes that translate to decorations on the result <id>
   LogicalResult processDecoration(Location loc, uint32_t resultID,
Serializer::processMemberDecoration(uint32_t structID, uint32_t memberIndex,
   if (!values.empty()) {
     args.append(values.begin(), values.end());
   }
   return encodeInstructionInto(decorations, spirv::Opcode::OpMemberDecorate,
                                args);
 }
 } // namespace
 
-LogicalResult Serializer::processFuncOp(FuncOp op) {
+LogicalResult Serializer::processFuncOp(spirv::FuncOp op) {
   LLVM_DEBUG(llvm::dbgs() << "-- start function '" << op.getName() << "' --\n");
   assert(functionHeader.empty() && functionBody.empty());
 
   uint32_t fnTypeID = 0;
   // Generate type of the function.
   processType(op.getLoc(), op.getType(), fnTypeID);
 
   // Add the function definition.
LogicalResult Serializer::processOperation(Operation *opInst) {
   // the SPIR-V spec.
   return TypeSwitch<Operation *, LogicalResult>(opInst)
       .Case([&](spirv::AddressOfOp op) { return processAddressOfOp(op); })
       .Case([&](spirv::BranchOp op) { return processBranchOp(op); })
       .Case([&](spirv::BranchConditionalOp op) {
         return processBranchConditionalOp(op);
       })
       .Case([&](spirv::ConstantOp op) { return processConstantOp(op); })
-      .Case([&](FuncOp op) { return processFuncOp(op); })
+      .Case([&](spirv::FuncOp op) { return processFuncOp(op); })
       .Case([&](spirv::GlobalVariableOp op) {
         return processGlobalVariableOp(op);
       })
       .Case([&](spirv::LoopOp op) { return processLoopOp(op); })
       .Case([&](spirv::ModuleEndOp) { return success(); })
       .Case([&](spirv::ReferenceOfOp op) { return processReferenceOfOp(op); })
       .Case([&](spirv::SelectionOp op) { return processSelectionOp(op); })
       .Case([&](spirv::SpecConstantOp op) { return processSpecConstantOp(op); })

mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/LayoutUtils.h"
 #include "mlir/Dialect/SPIRV/Passes.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVLowering.h"
 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
-#include "mlir/Dialect/StandardOps/Ops.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "llvm/ADT/SetVector.h"
 
 using namespace mlir;
 
 /// Checks if the `type` is a scalar or vector type. It is assumed that they are
 /// valid for SPIR-V dialect already.
 static bool isScalarOrVectorType(Type type) {
   return spirv::SPIRVDialect::isValidScalarType(type) || type.isa<VectorType>();
 }
 
 /// Creates a global variable for an argument based on the ABI info.
 static spirv::GlobalVariableOp
-createGlobalVariableForArg(FuncOp funcOp, OpBuilder &builder, unsigned argNum,
+createGlobalVariableForArg(spirv::FuncOp funcOp, OpBuilder &builder,
+                           unsigned argNum,
                            spirv::InterfaceVarABIAttr abiInfo) {
   auto spirvModule = funcOp.getParentOfType<spirv::ModuleOp>();
   if (!spirvModule) {
     return nullptr;
   }
   OpBuilder::InsertionGuard moduleInsertionGuard(builder);
   builder.setInsertionPoint(funcOp.getOperation());
   std::string varName =
createGlobalVariableForArg(spirv::FuncOp funcOp, OpBuilder &builder,
   return builder.create<spirv::GlobalVariableOp>(
       funcOp.getLoc(), varType, varName, abiInfo.descriptor_set().getInt(),
       abiInfo.binding().getInt());
 }
 
 /// Gets the global variables that need to be specified as interface variable
 /// with an spv.EntryPointOp. Traverses the body of a entry function to do so.
 static LogicalResult
-getInterfaceVariables(FuncOp funcOp,
+getInterfaceVariables(spirv::FuncOp funcOp,
                       SmallVectorImpl<Attribute> &interfaceVars) {
   auto module = funcOp.getParentOfType<spirv::ModuleOp>();
   if (!module) {
     return failure();
   }
   llvm::SetVector<Operation *> interfaceVarSet;
 
   // TODO(ravishankarm) : This should in reality traverse the entry function
getInterfaceVariables(spirv::FuncOp funcOp,
   for (auto &var : interfaceVarSet) {
     interfaceVars.push_back(SymbolRefAttr::get(
         cast<spirv::GlobalVariableOp>(var).sym_name(), funcOp.getContext()));
   }
   return success();
 }
 
 /// Lowers the entry point attribute.
-static LogicalResult lowerEntryPointABIAttr(FuncOp funcOp, OpBuilder &builder) {
+static LogicalResult lowerEntryPointABIAttr(spirv::FuncOp funcOp,
+                                            OpBuilder &builder) {
   auto entryPointAttrName = spirv::getEntryPointABIAttrName();
   auto entryPointAttr =
       funcOp.getAttrOfType<spirv::EntryPointABIAttr>(entryPointAttrName);
   if (!entryPointAttr) {
     return failure();
   }
 
   OpBuilder::InsertionGuard moduleInsertionGuard(builder);
static LogicalResult lowerEntryPointABIAttr(spirv::FuncOp funcOp,
   SmallVector<int32_t, 3> localSize(localSizeAttr.getValues<int32_t>());
   builder.create<spirv::ExecutionModeOp>(
       funcOp.getLoc(), funcOp, spirv::ExecutionMode::LocalSize, localSize);
   funcOp.removeAttr(entryPointAttrName);
   return success();
 }
 
 namespace {
-/// Pattern rewriter for changing function signature to match the ABI specified
-/// in attributes.
-class FuncOpLowering final : public SPIRVOpLowering<FuncOp> {
+/// A pattern to convert function signature according to interface variable ABI
+/// attributes.
+///
+/// Specifically, this pattern creates global variables according to interface
+/// variable ABI attributes attached to function arguments and converts all
+/// function argument uses to those global variables. This is necessary because
+/// Vulkan requires all shader entry points to be of void(void) type.
+class ProcessInterfaceVarABI final : public SPIRVOpLowering<spirv::FuncOp> {
 public:
-  using SPIRVOpLowering<FuncOp>::SPIRVOpLowering;
+  using SPIRVOpLowering<spirv::FuncOp>::SPIRVOpLowering;
   PatternMatchResult
-  matchAndRewrite(FuncOp funcOp, ArrayRef<Value> operands,
+  matchAndRewrite(spirv::FuncOp funcOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Pass to implement the ABI information specified as attributes.
 class LowerABIAttributesPass final
     : public OperationPass<LowerABIAttributesPass, spirv::ModuleOp> {
 private:
   void runOnOperation() override;
 };
 } // namespace
 
-PatternMatchResult
-FuncOpLowering::matchAndRewrite(FuncOp funcOp, ArrayRef<Value> operands,
-                                ConversionPatternRewriter &rewriter) const {
+PatternMatchResult ProcessInterfaceVarABI::matchAndRewrite(
+    spirv::FuncOp funcOp, ArrayRef<Value> operands,
   if (!funcOp.getAttrOfType<spirv::EntryPointABIAttr>(
           spirv::getEntryPointABIAttrName())) {
     // TODO(ravishankarm) : Non-entry point functions are not handled.
     return matchFailure();
   }
   TypeConverter::SignatureConversion signatureConverter(
       funcOp.getType().getNumInputs());
 
Value replacement =
         rewriter.create<spirv::AddressOfOp>(funcOp.getLoc(), var);
     // Check if the arg is a scalar or vector type. In that case, the value
     // needs to be loaded into registers.
     // TODO(ravishankarm) : This is loading value of the scalar into registers
     // at the start of the function. It is probably better to do the load just
     // before the use. There might be multiple loads and currently there is no
     // easy way to replace all uses with a sequence of operations.
     if (isScalarOrVectorType(argType.value())) {
-      auto indexType =
-          typeConverter.convertType(IndexType::get(funcOp.getContext()));
+      auto indexType = SPIRVTypeConverter::getIndexType(funcOp.getContext());
       auto zero =
           spirv::ConstantOp::getZero(indexType, funcOp.getLoc(), &rewriter);
       auto loadPtr = rewriter.create<spirv::AccessChainOp>(
           funcOp.getLoc(), replacement, zero.constant());
       replacement = rewriter.create<spirv::LoadOp>(funcOp.getLoc(), loadPtr);
     }
     signatureConverter.remapInput(argType.index(), replacement);
   }

 void LowerABIAttributesPass::runOnOperation() {
   // Uses the signature conversion methodology of the dialect conversion
   // framework to implement the conversion.
   spirv::ModuleOp module = getOperation();
   MLIRContext *context = &getContext();
 
   SPIRVTypeConverter typeConverter;
   OwningRewritePatternList patterns;
-  patterns.insert<FuncOpLowering>(context, typeConverter);
+  patterns.insert<ProcessInterfaceVarABI>(context, typeConverter);
 
-  std::unique_ptr<ConversionTarget> target = spirv::SPIRVConversionTarget::get(
-      spirv::lookupTargetEnvOrDefault(module), context);
-  auto entryPointAttrName = spirv::getEntryPointABIAttrName();
-  target->addDynamicallyLegalOp<FuncOp>([&](FuncOp op) {
-    return op.getAttrOfType<spirv::EntryPointABIAttr>(entryPointAttrName) &&
-           op.getNumResults() == 0 && op.getNumArguments() == 0;
+  ConversionTarget target(*context);
+  // "Legal" function ops should have no interface variable ABI attributes.
+  target.addDynamicallyLegalOp<spirv::FuncOp>([&](spirv::FuncOp op) {
+    StringRef attrName = spirv::getInterfaceVarABIAttrName();
+    for (unsigned i = 0, e = op.getNumArguments(); i < e; ++i)
+      if (op.getArgAttr(i, attrName))
+        return false;
+    return true;
+  });
+  // All other SPIR-V ops are legal.
+  target.markUnknownOpDynamicallyLegal([](Operation *op) {
+    return op->getDialect()->getNamespace() ==
+           spirv::SPIRVDialect::getDialectNamespace();
   });
-  target->addLegalOp<ReturnOp>();
   if (failed(
-          applyPartialConversion(module, *target, patterns, &typeConverter))) {
+          applyPartialConversion(module, target, patterns, &typeConverter))) {
     return signalPassFailure();
   }
 
   // Walks over all the FuncOps in spirv::ModuleOp to lower the entry point
   // attributes.
   OpBuilder builder(context);
-  SmallVector<FuncOp, 1> entryPointFns;
-  module.walk([&](FuncOp funcOp) {
+  SmallVector<spirv::FuncOp, 1> entryPointFns;
+  auto entryPointAttrName = spirv::getEntryPointABIAttrName();
+  module.walk([&](spirv::FuncOp funcOp) {
     if (funcOp.getAttrOfType<spirv::EntryPointABIAttr>(entryPointAttrName)) {
       entryPointFns.push_back(funcOp);
     }
   });
   for (auto fn : entryPointFns) {
     if (failed(lowerEntryPointABIAttr(fn, builder))) {
       return signalPassFailure();
     }

mlir/lib/IR/Function.cpp

void FuncOp::eraseArguments(ArrayRef<unsigned> argIndices) {
   // We do this in reverse so that we erase later indices before earlier
   // indices, to avoid shifting the later indices.
   Block &entry = front();
   for (int i = 0; i < originalNumArgs; i++)
     if (shouldEraseArg(originalNumArgs - i - 1))
       entry.eraseArgument(originalNumArgs - i - 1);
 }
 
-/// Add an entry block to an empty function, and set up the block arguments
-/// to match the signature of the function.
-Block *FuncOp::addEntryBlock() {
-  assert(empty() && "function already has an entry block");
-  auto *entry = new Block();
-  push_back(entry);
-  entry->addArguments(getType().getInputs());
-  return entry;
-}
-
-/// Add a normal block to the end of the function's block list. The function
-/// should at least already have an entry block.
-Block *FuncOp::addBlock() {
-  assert(!empty() && "function should at least have an entry block");
-  push_back(new Block());
-  return &back();
-}
-
 /// Clone the internal blocks from this function into dest and all attributes
 /// from this function to dest.
 void FuncOp::cloneInto(FuncOp dest, BlockAndValueMapping &mapper) {
   // Add the attributes of this function to dest.
   llvm::MapVector<Identifier, Attribute> newAttrs;
   for (auto &attr : dest.getAttrs())
     newAttrs.insert(attr);
   for (auto &attr : getAttrs())

mlir/test/Conversion/GPUToSPIRV/load-store.mlir

func @load_store(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>) {
     return
   }
 
   // CHECK-LABEL: spv.module "Logical" "GLSL450"
   gpu.module @kernels {
     // CHECK-DAG: spv.globalVariable [[NUMWORKGROUPSVAR:@.*]] built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
     // CHECK-DAG: spv.globalVariable [[LOCALINVOCATIONIDVAR:@.*]] built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
     // CHECK-DAG: spv.globalVariable [[WORKGROUPIDVAR:@.*]] built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
-    // CHECK-LABEL:    func @load_store_kernel
+    // CHECK-LABEL:    spv.func @load_store_kernel
     // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: [[ARG2:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: [[ARG3:%.*]]: i32 {spv.interface_var_abi = {binding = 3 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: [[ARG4:%.*]]: i32 {spv.interface_var_abi = {binding = 4 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: [[ARG5:%.*]]: i32 {spv.interface_var_abi = {binding = 5 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: [[ARG6:%.*]]: i32 {spv.interface_var_abi = {binding = 6 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     gpu.func @load_store_kernel(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>, %arg3: index, %arg4: index, %arg5: index, %arg6: index)

mlir/test/Conversion/GPUToSPIRV/simple.mlir

 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // CHECK:       spv.module "Logical" "GLSL450" {
-    // CHECK-LABEL: func @basic_module_structure
+    // CHECK-LABEL: spv.func @basic_module_structure
     // CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: {{%.*}}: !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
     gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32>)
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // CHECK: spv.Return
       gpu.return
     }

mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir

V_1_3, [], [Shader, GroupNonUniformArithmetic],
       max_compute_workgroup_invocations = 128 : i32,
       max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
     }>
 } {
 
 // CHECK:      spv.globalVariable
 // CHECK-SAME: built_in("LocalInvocationId")
 
-// CHECK:      func @single_workgroup_reduction
+// CHECK:      @single_workgroup_reduction
 // CHECK-SAME: (%[[INPUT:.+]]: !spv.ptr{{.+}}, %[[OUTPUT:.+]]: !spv.ptr{{.+}})
 
 // CHECK:        %[[ZERO:.+]] = spv.constant 0 : i32
 // CHECK:        %[[ID:.+]] = spv.Load "Input" %{{.+}} : vector<3xi32>
 // CHECK:        %[[X:.+]] = spv.CompositeExtract %[[ID]][0 : i32]
 
 // CHECK:        %[[INPTR:.+]] = spv.AccessChain %[[INPUT]][%[[ZERO]], %[[X]]]
 // CHECK:        %[[VAL:.+]] = spv.Load "StorageBuffer" %[[INPTR]] : i32

mlir/test/Conversion/StandardToSPIRV/std-to-spirv.mlir

 // RUN: mlir-opt -convert-std-to-spirv %s -o - | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // std binary arithmetic ops
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @add_sub
 func @add_sub(%arg0 : i32, %arg1 : i32) {
   // CHECK: spv.IAdd
   %0 = addi %arg0, %arg1 : i32
   // CHECK: spv.ISub
   %1 = subi %arg0, %arg1 : i32
   return
 }
 
 // CHECK-LABEL: @fadd_scalar
-func @fadd_scalar(%arg: f32) -> f32 {
+func @fadd_scalar(%arg: f32) {
   // CHECK: spv.FAdd
   %0 = addf %arg, %arg : f32
-  return %0 : f32
+  return
 }
 
 // CHECK-LABEL: @fdiv_scalar
-func @fdiv_scalar(%arg: f32) -> f32 {
+func @fdiv_scalar(%arg: f32) {
   // CHECK: spv.FDiv
   %0 = divf %arg, %arg : f32
-  return %0 : f32
+  return
 }
 
 // CHECK-LABEL: @fmul_scalar
-func @fmul_scalar(%arg: f32) -> f32 {
+func @fmul_scalar(%arg: f32) {
   // CHECK: spv.FMul
   %0 = mulf %arg, %arg : f32
-  return %0 : f32
+  return
 }
 
 // CHECK-LABEL: @fmul_vector2
-func @fmul_vector2(%arg: vector<2xf32>) -> vector<2xf32> {
+func @fmul_vector2(%arg: vector<2xf32>) {
   // CHECK: spv.FMul
   %0 = mulf %arg, %arg : vector<2xf32>
-  return %0 : vector<2xf32>
+  return
 }
 
 // CHECK-LABEL: @fmul_vector3
-func @fmul_vector3(%arg: vector<3xf32>) -> vector<3xf32> {
+func @fmul_vector3(%arg: vector<3xf32>) {
   // CHECK: spv.FMul
   %0 = mulf %arg, %arg : vector<3xf32>
-  return %0 : vector<3xf32>
+  return
 }
 
 // CHECK-LABEL: @fmul_vector4
-func @fmul_vector4(%arg: vector<4xf32>) -> vector<4xf32> {
+func @fmul_vector4(%arg: vector<4xf32>) {
   // CHECK: spv.FMul
   %0 = mulf %arg, %arg : vector<4xf32>
-  return %0 : vector<4xf32>
+  return
 }
 
 // CHECK-LABEL: @fmul_vector5
-func @fmul_vector5(%arg: vector<5xf32>) -> vector<5xf32> {
-  // Vector length of only 2, 3, and 4 is valid for SPIR-V
+func @fmul_vector5(%arg: vector<5xf32>) {
+  // Vector length of only 2, 3, and 4 is valid for SPIR-V.
   // CHECK: mulf
   %0 = mulf %arg, %arg : vector<5xf32>
-  return %0 : vector<5xf32>
+  return
 }
 
-// CHECK-LABEL: @fmul_tensor
-func @fmul_tensor(%arg: tensor<4xf32>) -> tensor<4xf32> {
-  // For tensors mulf cannot be lowered directly to spv.FMul
-  // CHECK: mulf
-  %0 = mulf %arg, %arg : tensor<4xf32>
-  return %0 : tensor<4xf32>
-}
+// TODO(antiagainst): enable this once we support converting binary ops
+// needing type conversion.
+// XXXXX-LABEL: @fmul_tensor
+//func @fmul_tensor(%arg: tensor<4xf32>) {
+  // For tensors mulf cannot be lowered directly to spv.FMul.
+  // XXXXX: mulf
+  //%0 = mulf %arg, %arg : tensor<4xf32>
+  //return
+//}
 
 // CHECK-LABEL: @frem_scalar
-func @frem_scalar(%arg: f32) -> f32 {
+func @frem_scalar(%arg: f32) {
   // CHECK: spv.FRem
   %0 = remf %arg, %arg : f32
-  return %0 : f32
+  return
 }
 
 // CHECK-LABEL: @fsub_scalar
-func @fsub_scalar(%arg: f32) -> f32 {
+func @fsub_scalar(%arg: f32) {
   // CHECK: spv.FSub
   %0 = subf %arg, %arg : f32
-  return %0 : f32
+  return
 }
 
 // CHECK-LABEL: @div_rem
 func @div_rem(%arg0 : i32, %arg1 : i32) {
   // CHECK: spv.SDiv
   %0 = divi_signed %arg0, %arg1 : i32
   // CHECK: spv.SMod
   %1 = remi_signed %arg0, %arg1 : i32
func @sitofp(%arg0 : i32) {
   %0 = std.sitofp %arg0 : i32 to f32
   return
 }
 
 //===----------------------------------------------------------------------===//
 // memref type
 //===----------------------------------------------------------------------===//
 
-// CHECK-LABEL: func @memref_type({{%.*}}: memref<3xi1>) {
+// CHECK-LABEL: func @memref_type({{%.*}}: memref<3xi1>)
 func @memref_type(%arg0: memref<3xi1>) {
   return
 }

mlir/test/Dialect/SPIRV/Serialization/arithmetic-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @fmul(%arg0 : f32, %arg1 : f32) {
+  spv.func @fmul(%arg0 : f32, %arg1 : f32) "None" {
     // CHECK: {{%.*}}= spv.FMul {{%.*}}, {{%.*}} : f32
     %0 = spv.FMul %arg0, %arg1 : f32
     spv.Return
   }
-  func @fadd(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) {
+  spv.func @fadd(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: {{%.*}} = spv.FAdd {{%.*}}, {{%.*}} : vector<4xf32>
     %0 = spv.FAdd %arg0, %arg1 : vector<4xf32>
     spv.Return
   }
-  func @fdiv(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) {
+  spv.func @fdiv(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: {{%.*}} = spv.FDiv {{%.*}}, {{%.*}} : vector<4xf32>
     %0 = spv.FDiv %arg0, %arg1 : vector<4xf32>
     spv.Return
   }
-  func @fmod(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) {
+  spv.func @fmod(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: {{%.*}} = spv.FMod {{%.*}}, {{%.*}} : vector<4xf32>
     %0 = spv.FMod %arg0, %arg1 : vector<4xf32>
     spv.Return
   }
-  func @fnegate(%arg0 : vector<4xf32>) {
+  spv.func @fnegate(%arg0 : vector<4xf32>) "None" {
     // CHECK: {{%.*}} = spv.FNegate {{%.*}} : vector<4xf32>
     %0 = spv.FNegate %arg0 : vector<4xf32>
     spv.Return
   }
-  func @fsub(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) {
+  spv.func @fsub(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: {{%.*}} = spv.FSub {{%.*}}, {{%.*}} : vector<4xf32>
     %0 = spv.FSub %arg0, %arg1 : vector<4xf32>
     spv.Return
   }
-  func @frem(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) {
+  spv.func @frem(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: {{%.*}} = spv.FRem {{%.*}}, {{%.*}} : vector<4xf32>
     %0 = spv.FRem %arg0, %arg1 : vector<4xf32>
     spv.Return
   }
-  func @iadd(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
+  spv.func @iadd(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) "None" {
     // CHECK: {{%.*}} = spv.IAdd {{%.*}}, {{%.*}} : vector<4xi32>
     %0 = spv.IAdd %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @isub(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
+  spv.func @isub(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) "None" {
     // CHECK: {{%.*}} = spv.ISub {{%.*}}, {{%.*}} : vector<4xi32>
     %0 = spv.ISub %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @imul(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
+  spv.func @imul(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) "None" {
     // CHECK: {{%.*}} = spv.IMul {{%.*}}, {{%.*}} : vector<4xi32>
     %0 = spv.IMul %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @udiv(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
+  spv.func @udiv(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) "None" {
     // CHECK: {{%.*}} = spv.UDiv {{%.*}}, {{%.*}} : vector<4xi32>
     %0 = spv.UDiv %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @umod(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
+  spv.func @umod(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) "None" {
     // CHECK: {{%.*}} = spv.UMod {{%.*}}, {{%.*}} : vector<4xi32>
     %0 = spv.UMod %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @sdiv(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
+  spv.func @sdiv(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) "None" {
     // CHECK: {{%.*}} = spv.SDiv {{%.*}}, {{%.*}} : vector<4xi32>
     %0 = spv.SDiv %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @smod(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
+  spv.func @smod(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) "None" {
     // CHECK: {{%.*}} = spv.SMod {{%.*}}, {{%.*}} : vector<4xi32>
     %0 = spv.SMod %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @srem(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) {
+  spv.func @srem(%arg0 : vector<4xi32>, %arg1 : vector<4xi32>) "None" {
     // CHECK: {{%.*}} = spv.SRem {{%.*}}, {{%.*}} : vector<4xi32>
     %0 = spv.SRem %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/array.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @array_stride(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32 [4]> [128]>, StorageBuffer>,
-                     %arg1 : i32, %arg2 : i32) {
+  spv.func @array_stride(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32 [4]> [128]>, StorageBuffer>, %arg1 : i32, %arg2 : i32) "None" {
     // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32 [4]> [128]>, StorageBuffer>
     %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32 [4]> [128]>, StorageBuffer>
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.rtarray<f32>, StorageBuffer>
   spv.globalVariable @var0 : !spv.ptr<!spv.rtarray<f32>, StorageBuffer>
   // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
   spv.globalVariable @var1 : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
 }

mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
   // CHECK-LABEL: @atomic_compare_exchange_weak
-  func @atomic_compare_exchange_weak(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 {
+  spv.func @atomic_compare_exchange_weak(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 "None" {
     // CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %{{.*}}, %{{.*}}, %{{.*}} : !spv.ptr<i32, Workgroup>
     %0 = spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %ptr, %value, %comparator: !spv.ptr<i32, Workgroup>
     // CHECK: spv.AtomicAnd "Device" "None" %{{.*}}, %{{.*}} : !spv.ptr<i32, Workgroup>
     %1 = spv.AtomicAnd "Device" "None" %ptr, %value : !spv.ptr<i32, Workgroup>
     // CHECK: spv.AtomicIAdd "Workgroup" "Acquire" %{{.*}}, %{{.*}} : !spv.ptr<i32, Workgroup>
     %2 = spv.AtomicIAdd "Workgroup" "Acquire" %ptr, %value : !spv.ptr<i32, Workgroup>
     // CHECK: spv.AtomicIDecrement "Workgroup" "Acquire" %{{.*}} : !spv.ptr<i32, Workgroup>
     %3 = spv.AtomicIDecrement "Workgroup" "Acquire" %ptr : !spv.ptr<i32, Workgroup>

mlir/test/Dialect/SPIRV/Serialization/barrier.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @memory_barrier_0() -> () {
+  spv.func @memory_barrier_0() -> () "None" {
     // CHECK: spv.MemoryBarrier "Device", "Release|UniformMemory"
     spv.MemoryBarrier "Device", "Release|UniformMemory"
     spv.Return
   }
-  func @memory_barrier_1() -> () {
+  spv.func @memory_barrier_1() -> () "None" {
     // CHECK: spv.MemoryBarrier "Subgroup", "AcquireRelease|SubgroupMemory"
     spv.MemoryBarrier "Subgroup", "AcquireRelease|SubgroupMemory"
     spv.Return
   }
-  func @control_barrier_0() -> () {
+  spv.func @control_barrier_0() -> () "None" {
     // CHECK: spv.ControlBarrier "Device", "Workgroup", "Release|UniformMemory"
     spv.ControlBarrier "Device", "Workgroup", "Release|UniformMemory"
     spv.Return
   }
-  func @control_barrier_1() -> () {
+  spv.func @control_barrier_1() -> () "None" {
     // CHECK: spv.ControlBarrier "Workgroup", "Invocation", "AcquireRelease|UniformMemory"
     spv.ControlBarrier "Workgroup", "Invocation", "AcquireRelease|UniformMemory"
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/bit-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @bitcount(%arg: i32) -> i32 {
+  spv.func @bitcount(%arg: i32) -> i32 "None" {
     // CHECK: spv.BitCount {{%.*}} : i32
     %0 = spv.BitCount %arg : i32
     spv.ReturnValue %0 : i32
   }
-  func @bit_field_insert(%base: vector<3xi32>, %insert: vector<3xi32>, %offset: i32, %count: i16) -> vector<3xi32> {
+  spv.func @bit_field_insert(%base: vector<3xi32>, %insert: vector<3xi32>, %offset: i32, %count: i16) -> vector<3xi32> "None" {
     // CHECK: {{%.*}} = spv.BitFieldInsert {{%.*}}, {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi32>, i32, i16
     %0 = spv.BitFieldInsert %base, %insert, %offset, %count : vector<3xi32>, i32, i16
     spv.ReturnValue %0 : vector<3xi32>
   }
-  func @bit_field_s_extract(%base: vector<3xi32>, %offset: i8, %count: i8) -> vector<3xi32> {
+  spv.func @bit_field_s_extract(%base: vector<3xi32>, %offset: i8, %count: i8) -> vector<3xi32> "None" {
     // CHECK: {{%.*}} = spv.BitFieldSExtract {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi32>, i8, i8
     %0 = spv.BitFieldSExtract %base, %offset, %count : vector<3xi32>, i8, i8
     spv.ReturnValue %0 : vector<3xi32>
   }
-  func @bit_field_u_extract(%base: vector<3xi32>, %offset: i8, %count: i8) -> vector<3xi32> {
+  spv.func @bit_field_u_extract(%base: vector<3xi32>, %offset: i8, %count: i8) -> vector<3xi32> "None" {
     // CHECK: {{%.*}} = spv.BitFieldUExtract {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi32>, i8, i8
     %0 = spv.BitFieldUExtract %base, %offset, %count : vector<3xi32>, i8, i8
     spv.ReturnValue %0 : vector<3xi32>
   }
-  func @bitreverse(%arg: i32) -> i32 {
+  spv.func @bitreverse(%arg: i32) -> i32 "None" {
     // CHECK: spv.BitReverse {{%.*}} : i32
     %0 = spv.BitReverse %arg : i32
     spv.ReturnValue %0 : i32
   }
-  func @not(%arg: i32) -> i32 {
+  spv.func @not(%arg: i32) -> i32 "None" {
     // CHECK: spv.Not {{%.*}} : i32
     %0 = spv.Not %arg : i32
     spv.ReturnValue %0 : i32
   }
-  func @bitwise_scalar(%arg0 : i32, %arg1 : i32) {
+  spv.func @bitwise_scalar(%arg0 : i32, %arg1 : i32) "None" {
     // CHECK: spv.BitwiseAnd
     %0 = spv.BitwiseAnd %arg0, %arg1 : i32
     // CHECK: spv.BitwiseOr
     %1 = spv.BitwiseOr %arg0, %arg1 : i32
     // CHECK: spv.BitwiseXor
     %2 = spv.BitwiseXor %arg0, %arg1 : i32
     spv.Return
   }
-  func @shift_left_logical(%arg0: i32, %arg1 : i16) -> i32 {
+  spv.func @shift_left_logical(%arg0: i32, %arg1 : i16) -> i32 "None" {
     // CHECK: {{%.*}} = spv.ShiftLeftLogical {{%.*}}, {{%.*}} : i32, i16
     %0 = spv.ShiftLeftLogical %arg0, %arg1: i32, i16
     spv.ReturnValue %0 : i32
   }
-  func @shift_right_arithmetic(%arg0: vector<4xi32>, %arg1 : vector<4xi8>) -> vector<4xi32> {
+  spv.func @shift_right_arithmetic(%arg0: vector<4xi32>, %arg1 : vector<4xi8>) -> vector<4xi32> "None" {
     // CHECK: {{%.*}} = spv.ShiftRightArithmetic {{%.*}}, {{%.*}} : vector<4xi32>, vector<4xi8>
     %0 = spv.ShiftRightArithmetic %arg0, %arg1: vector<4xi32>, vector<4xi8>
     spv.ReturnValue %0 : vector<4xi32>
   }
-  func @shift_right_logical(%arg0: vector<2xi32>, %arg1 : vector<2xi8>) -> vector<2xi32> {
+  spv.func @shift_right_logical(%arg0: vector<2xi32>, %arg1 : vector<2xi8>) -> vector<2xi32> "None" {
     // CHECK: {{%.*}} = spv.ShiftRightLogical {{%.*}}, {{%.*}} : vector<2xi32>, vector<2xi8>
     %0 = spv.ShiftRightLogical %arg0, %arg1: vector<2xi32>, vector<2xi8>
     spv.ReturnValue %0 : vector<2xi32>
   }
 }

mlir/test/Dialect/SPIRV/Serialization/cast-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @bit_cast(%arg0 : f32) {
+  spv.func @bit_cast(%arg0 : f32) "None" {
     // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : f32 to i32
     %0 = spv.Bitcast %arg0 : f32 to i32
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @convert_f_to_s(%arg0 : f32) -> i32 {
+  spv.func @convert_f_to_s(%arg0 : f32) -> i32 "None" {
     // CHECK: {{%.*}} = spv.ConvertFToS {{%.*}} : f32 to i32
     %0 = spv.ConvertFToS %arg0 : f32 to i32
     spv.ReturnValue %0 : i32
   }
-  func @convert_f_to_u(%arg0 : f32) -> i32 {
+  spv.func @convert_f_to_u(%arg0 : f32) -> i32 "None" {
     // CHECK: {{%.*}} = spv.ConvertFToU {{%.*}} : f32 to i32
     %0 = spv.ConvertFToU %arg0 : f32 to i32
     spv.ReturnValue %0 : i32
   }
-  func @convert_s_to_f(%arg0 : i32) -> f32 {
+  spv.func @convert_s_to_f(%arg0 : i32) -> f32 "None" {
     // CHECK: {{%.*}} = spv.ConvertSToF {{%.*}} : i32 to f32
     %0 = spv.ConvertSToF %arg0 : i32 to f32
     spv.ReturnValue %0 : f32
   }
-  func @convert_u_to_f(%arg0 : i32) -> f32 {
+  spv.func @convert_u_to_f(%arg0 : i32) -> f32 "None" {
     // CHECK: {{%.*}} = spv.ConvertUToF {{%.*}} : i32 to f32
     %0 = spv.ConvertUToF %arg0 : i32 to f32
     spv.ReturnValue %0 : f32
   }
-  func @f_convert(%arg0 : f32) -> f64 {
+  spv.func @f_convert(%arg0 : f32) -> f64 "None" {
     // CHECK: {{%.*}} = spv.FConvert {{%.*}} : f32 to f64
     %0 = spv.FConvert %arg0 : f32 to f64
     spv.ReturnValue %0 : f64
   }
-  func @s_convert(%arg0 : i32) -> i64 {
+  spv.func @s_convert(%arg0 : i32) -> i64 "None" {
     // CHECK: {{%.*}} = spv.SConvert {{%.*}} : i32 to i64
     %0 = spv.SConvert %arg0 : i32 to i64
     spv.ReturnValue %0 : i64
   }
-  func @u_convert(%arg0 : i32) -> i64 {
+  spv.func @u_convert(%arg0 : i32) -> i64 "None" {
     // CHECK: {{%.*}} = spv.UConvert {{%.*}} : i32 to i64
     %0 = spv.UConvert %arg0 : i32 to i64
     spv.ReturnValue %0 : i64
   }
 }

mlir/test/Dialect/SPIRV/Serialization/composite-op.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @composite_insert(%arg0 : !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>, %arg1: !spv.array<4xf32>) -> !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>> {
+  spv.func @composite_insert(%arg0 : !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>, %arg1: !spv.array<4xf32>) -> !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>> "None" {
     // CHECK: spv.CompositeInsert {{%.*}}, {{%.*}}[1 : i32, 0 : i32] : !spv.array<4 x f32> into !spv.struct<f32, !spv.struct<!spv.array<4 x f32>, f32>>
     %0 = spv.CompositeInsert %arg1, %arg0[1 : i32, 0 : i32] : !spv.array<4xf32> into !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>
     spv.ReturnValue %0: !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>
   }
-  func @composite_construct_vector(%arg0: f32, %arg1: f32, %arg2 : f32) -> vector<3xf32> {
+  spv.func @composite_construct_vector(%arg0: f32, %arg1: f32, %arg2 : f32) -> vector<3xf32> "None" {
     // CHECK: spv.CompositeConstruct {{%.*}}, {{%.*}}, {{%.*}} : vector<3xf32>
     %0 = spv.CompositeConstruct %arg0, %arg1, %arg2 : vector<3xf32>
     spv.ReturnValue %0: vector<3xf32>
   }
 }

mlir/test/Dialect/SPIRV/Serialization/constant.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
   // CHECK-LABEL: @bool_const
-  func @bool_const() -> () {
+  spv.func @bool_const() -> () "None" {
     // CHECK: spv.constant true
     %0 = spv.constant true
     // CHECK: spv.constant false
     %1 = spv.constant false
 
     %2 = spv.Variable init(%0): !spv.ptr<i1, Function>
     %3 = spv.Variable init(%1): !spv.ptr<i1, Function>
     spv.Return
   }
 
   // CHECK-LABEL: @i32_const
-  func @i32_const() -> () {
+  spv.func @i32_const() -> () "None" {
     // CHECK: spv.constant 0 : i32
     %0 = spv.constant  0 : i32
     // CHECK: spv.constant 10 : i32
     %1 = spv.constant 10 : i32
     // CHECK: spv.constant -5 : i32
     %2 = spv.constant -5 : i32
 
     %3 = spv.IAdd %0, %1 : i32
     %4 = spv.IAdd %2, %3 : i32
     spv.Return
   }
 
   // CHECK-LABEL: @i64_const
-  func @i64_const() -> () {
+  spv.func @i64_const() -> () "None" {
     // CHECK: spv.constant 4294967296 : i64
     %0 = spv.constant           4294967296 : i64 //  2^32
     // CHECK: spv.constant -4294967296 : i64
     %1 = spv.constant          -4294967296 : i64 // -2^32
     // CHECK: spv.constant 9223372036854775807 : i64
     %2 = spv.constant  9223372036854775807 : i64 //  2^63 - 1
     // CHECK: spv.constant -9223372036854775808 : i64
     %3 = spv.constant -9223372036854775808 : i64 // -2^63
 
     %4 = spv.IAdd %0, %1 : i64
     %5 = spv.IAdd %2, %3 : i64
     spv.Return
   }
 
   // CHECK-LABEL: @i16_const
-  func @i16_const() -> () {
+  spv.func @i16_const() -> () "None" {
     // CHECK: spv.constant -32768 : i16
     %0 = spv.constant -32768 : i16 // -2^15
     // CHECK: spv.constant 32767 : i16
     %1 = spv.constant 32767 : i16 //  2^15 - 1
 
     %2 = spv.IAdd %0, %1 : i16
     spv.Return
   }
 
   // CHECK-LABEL: @float_const
-  func @float_const() -> () {
+  spv.func @float_const() -> () "None" {
     // CHECK: spv.constant 0.000000e+00 : f32
     %0 = spv.constant 0. : f32
     // CHECK: spv.constant 1.000000e+00 : f32
     %1 = spv.constant 1. : f32
     // CHECK: spv.constant -0.000000e+00 : f32
     %2 = spv.constant -0. : f32
     // CHECK: spv.constant -1.000000e+00 : f32
     %3 = spv.constant -1. : f32
     // CHECK: spv.constant 7.500000e-01 : f32
     %4 = spv.constant 0.75 : f32
     // CHECK: spv.constant -2.500000e-01 : f32
     %5 = spv.constant -0.25 : f32
 
     %6 = spv.FAdd %0, %1 : f32
     %7 = spv.FAdd %2, %3 : f32
     %8 = spv.FAdd %4, %5 : f32
     spv.Return
   }
 
   // CHECK-LABEL: @double_const
-  func @double_const() -> () {
+  spv.func @double_const() -> () "None" {
     // TODO(antiagainst): test range boundary values
     // CHECK: spv.constant 1.024000e+03 : f64
     %0 = spv.constant 1024. : f64
     // CHECK: spv.constant -1.024000e+03 : f64
     %1 = spv.constant -1024. : f64
 
     %2 = spv.FAdd %0, %1 : f64
     spv.Return
   }
 
   // CHECK-LABEL: @half_const
-  func @half_const() -> () {
+  spv.func @half_const() -> () "None" {
     // CHECK: spv.constant 5.120000e+02 : f16
     %0 = spv.constant 512. : f16
     // CHECK: spv.constant -5.120000e+02 : f16
     %1 = spv.constant -512. : f16
 
     %2 = spv.FAdd %0, %1 : f16
     spv.Return
   }
 
   // CHECK-LABEL: @bool_vector_const
-  func @bool_vector_const() -> () {
+  spv.func @bool_vector_const() -> () "None" {
     // CHECK: spv.constant dense<false> : vector<2xi1>
     %0 = spv.constant dense<false> : vector<2xi1>
     // CHECK: spv.constant dense<[true, true, true]> : vector<3xi1>
     %1 = spv.constant dense<true> : vector<3xi1>
     // CHECK: spv.constant dense<[false, true]> : vector<2xi1>
     %2 = spv.constant dense<[false, true]> : vector<2xi1>
 
     %3 = spv.Variable init(%0): !spv.ptr<vector<2xi1>, Function>
     %4 = spv.Variable init(%1): !spv.ptr<vector<3xi1>, Function>
     %5 = spv.Variable init(%2): !spv.ptr<vector<2xi1>, Function>
     spv.Return
   }
 
   // CHECK-LABEL: @int_vector_const
-  func @int_vector_const() -> () {
+  spv.func @int_vector_const() -> () "None" {
     // CHECK: spv.constant dense<0> : vector<3xi32>
     %0 = spv.constant dense<0> : vector<3xi32>
     // CHECK: spv.constant dense<1> : vector<3xi32>
     %1 = spv.constant dense<1> : vector<3xi32>
     // CHECK: spv.constant dense<[2, -3, 4]> : vector<3xi32>
     %2 = spv.constant dense<[2, -3, 4]> : vector<3xi32>
 
     %3 = spv.IAdd %0, %1 : vector<3xi32>
     %4 = spv.IAdd %2, %3 : vector<3xi32>
     spv.Return
   }
 
   // CHECK-LABEL: @fp_vector_const
-  func @fp_vector_const() -> () {
+  spv.func @fp_vector_const() -> () "None" {
     // CHECK: spv.constant dense<0.000000e+00> : vector<4xf32>
     %0 = spv.constant dense<0.> : vector<4xf32>
     // CHECK: spv.constant dense<-1.500000e+01> : vector<4xf32>
     %1 = spv.constant dense<-15.> : vector<4xf32>
     // CHECK: spv.constant dense<[7.500000e-01, -2.500000e-01, 1.000000e+01, 4.200000e+01]> : vector<4xf32>
     %2 = spv.constant dense<[0.75, -0.25, 10., 42.]> : vector<4xf32>
 
     %3 = spv.FAdd %0, %1 : vector<4xf32>
     %4 = spv.FAdd %2, %3 : vector<4xf32>
     spv.Return
   }
 
   // CHECK-LABEL: @array_const
-  func @array_const() -> (!spv.array<2 x vector<2xf32>>) {
+  spv.func @array_const() -> (!spv.array<2 x vector<2xf32>>) "None" {
     // CHECK: spv.constant [dense<3.000000e+00> : vector<2xf32>, dense<[4.000000e+00, 5.000000e+00]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
     %0 = spv.constant [dense<3.0> : vector<2xf32>, dense<[4., 5.]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
 
     spv.ReturnValue %0 : !spv.array<2 x vector<2xf32>>
   }
 
   // CHECK-LABEL: @ignore_not_used_const
-  func @ignore_not_used_const() -> () {
+  spv.func @ignore_not_used_const() -> () "None" {
     %0 = spv.constant false
     // CHECK-NEXT: spv.Return
     spv.Return
   }
 
   // CHECK-LABEL: @materialize_const_at_each_use
-  func @materialize_const_at_each_use() -> (i32) {
+  spv.func @materialize_const_at_each_use() -> (i32) "None" {
     // CHECK: %[[USE1:.*]] = spv.constant 42 : i32
     // CHECK: %[[USE2:.*]] = spv.constant 42 : i32
     // CHECK: spv.IAdd %[[USE1]], %[[USE2]]
     %0 = spv.constant 42 : i32
     %1 = spv.IAdd %0, %0 : i32
     spv.ReturnValue %1 : i32
   }
 
   // CHECK-LABEL: @const_variable
-  func @const_variable(%arg0 : i32, %arg1 : i32) -> () {
+  spv.func @const_variable(%arg0 : i32, %arg1 : i32) -> () "None" {
     // CHECK: %[[CONST:.*]] = spv.constant 5 : i32
     // CHECK: spv.Variable init(%[[CONST]]) : !spv.ptr<i32, Function>
     // CHECK: spv.IAdd %arg0, %arg1
     %0 = spv.IAdd %arg0, %arg1 : i32
     %1 = spv.constant 5 : i32
     %2 = spv.Variable init(%1) : !spv.ptr<i32, Function>
     %3 = spv.Load "Function" %2 : i32
     %4 = spv.IAdd %0, %3 : i32
     spv.Return
   }
 
   // CHECK-LABEL: @multi_dimensions_const
-  func @multi_dimensions_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>) {
+  spv.func @multi_dimensions_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>) "None" {
     // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
     %0 = spv.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
     spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
   }
 
   // CHECK-LABEL: @multi_dimensions_splat_const
-  func @multi_dimensions_splat_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>) {
+  spv.func @multi_dimensions_splat_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>) "None" {
     // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
     %0 = spv.constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
     spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
   }
 }

mlir/test/Dialect/SPIRV/Serialization/entry-point.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @noop() -> () {
+  spv.func @noop() -> () "None" {
     spv.Return
   }
   // CHECK:      spv.EntryPoint "GLCompute" @noop
   // CHECK-NEXT: spv.ExecutionMode @noop "ContractionOff"
   spv.EntryPoint "GLCompute" @noop
   spv.ExecutionMode @noop "ContractionOff"
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   // CHECK:       spv.globalVariable @var2 : !spv.ptr<f32, Input>
   // CHECK-NEXT:  spv.globalVariable @var3 : !spv.ptr<f32, Output>
-  // CHECK-NEXT:  func @noop({{%.*}}: !spv.ptr<f32, Input>, {{%.*}}: !spv.ptr<f32, Output>)
+  // CHECK-NEXT:  spv.func @noop({{%.*}}: !spv.ptr<f32, Input>, {{%.*}}: !spv.ptr<f32, Output>) "None"
   // CHECK:       spv.EntryPoint "GLCompute" @noop, @var2, @var3
   spv.globalVariable @var2 : !spv.ptr<f32, Input>
   spv.globalVariable @var3 : !spv.ptr<f32, Output>
-  func @noop(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () {
+  spv.func @noop(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @noop, @var2, @var3
   spv.ExecutionMode @noop "ContractionOff"
 }

mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @foo
   // CHECK: spv.ExecutionMode @foo "LocalSizeHint", 3, 4, 5
   spv.ExecutionMode @foo "LocalSizeHint", 3, 4, 5
 }

mlir/test/Dialect/SPIRV/Serialization/function-call.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @var1 : !spv.ptr<!spv.array<4xf32>, Input>
-  func @fmain() -> i32 {
+  spv.func @fmain() -> i32 "None" {
     %0 = spv.constant 16 : i32
     %1 = spv._address_of @var1 : !spv.ptr<!spv.array<4xf32>, Input>
     // CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}) : (i32) -> i32
     %3 = spv.FunctionCall @f_0(%0) : (i32) -> i32
     // CHECK: spv.FunctionCall @f_1({{%.*}}, {{%.*}}) : (i32, !spv.ptr<!spv.array<4 x f32>, Input>) -> ()
     spv.FunctionCall @f_1(%3, %1) : (i32, !spv.ptr<!spv.array<4xf32>, Input>) ->  ()
     // CHECK: {{%.*}} =  spv.FunctionCall @f_2({{%.*}}) : (!spv.ptr<!spv.array<4 x f32>, Input>) -> !spv.ptr<!spv.array<4 x f32>, Input>
     %4 = spv.FunctionCall @f_2(%1) : (!spv.ptr<!spv.array<4xf32>, Input>) -> !spv.ptr<!spv.array<4xf32>, Input>
     spv.ReturnValue %3 : i32
   }
-  func @f_0(%arg0 : i32) -> i32 {
+  spv.func @f_0(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
-  func @f_1(%arg0 : i32, %arg1 : !spv.ptr<!spv.array<4xf32>, Input>) -> () {
+  spv.func @f_1(%arg0 : i32, %arg1 : !spv.ptr<!spv.array<4xf32>, Input>) -> () "None" {
     spv.Return
   }
-  func @f_2(%arg0 : !spv.ptr<!spv.array<4xf32>, Input>) -> !spv.ptr<!spv.array<4xf32>, Input> {
+  spv.func @f_2(%arg0 : !spv.ptr<!spv.array<4xf32>, Input>) -> !spv.ptr<!spv.array<4xf32>, Input> "None" {
     spv.ReturnValue %arg0 : !spv.ptr<!spv.array<4xf32>, Input>
   }
 
-  func @f_loop_with_function_call(%count : i32) -> () {
+  spv.func @f_loop_with_function_call(%count : i32) -> () "None" {
     %zero = spv.constant 0: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     spv.loop {
       spv.Branch ^header
     ^header:
       %val0 = spv.Load "Function" %var : i32
       %cmp = spv.SLessThan %val0, %count : i32
       spv.BranchConditional %cmp, ^body, ^merge
     ^body:
       spv.Branch ^continue
     ^continue:
       // CHECK: spv.FunctionCall @f_inc({{%.*}}) : (!spv.ptr<i32, Function>) -> ()
       spv.FunctionCall @f_inc(%var) : (!spv.ptr<i32, Function>) -> ()
       spv.Branch ^header
     ^merge:
       spv._merge
     }
     spv.Return
   }
-  func @f_inc(%arg0 : !spv.ptr<i32, Function>) -> () {
+  spv.func @f_inc(%arg0 : !spv.ptr<i32, Function>) -> () "None" {
       %one = spv.constant 1 : i32
       %0 = spv.Load "Function" %arg0 : i32
       %1 = spv.IAdd %0, %one : i32
       spv.Store "Function" %arg0, %1 : i32
       spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir

spv.module "Logical" "GLSL450" {
   spv.globalVariable @var1 : !spv.ptr<f32, Input>
   spv.globalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
-  func @foo() {
+  spv.func @foo() "None" {
     // CHECK: %[[ADDR:.*]] = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
     %0 = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
     %1 = spv.constant 0: i32
     // CHECK: spv.AccessChain %[[ADDR]]
     %2 = spv.AccessChain %0[%1] : !spv.ptr<vector<3xi32>, Input>
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @fmul(%arg0 : f32, %arg1 : f32) {
+  spv.func @fmul(%arg0 : f32, %arg1 : f32) "None" {
     // CHECK: {{%.*}} = spv.GLSL.Exp {{%.*}} : f32
     %0 = spv.GLSL.Exp %arg0 : f32
     // CHECK: {{%.*}} = spv.GLSL.FMax {{%.*}}, {{%.*}} : f32
     %1 = spv.GLSL.FMax %arg0, %arg1 : f32
     // CHECK: {{%.*}} = spv.GLSL.Sqrt {{%.*}} : f32
     %2 = spv.GLSL.Sqrt %arg0 : f32
     // CHECK: {{%.*}} = spv.GLSL.Cos {{%.*}} : f32
     %3 = spv.GLSL.Cos %arg0 : f32
     // CHECK: {{%.*}} = spv.GLSL.Sin {{%.*}} : f32
     %4 = spv.GLSL.Sin %arg0 : f32
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/group-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
   // CHECK-LABEL: @subgroup_ballot
-  func @subgroup_ballot(%predicate: i1) -> vector<4xi32> {
+  spv.func @subgroup_ballot(%predicate: i1) -> vector<4xi32> "None" {
     // CHECK: %{{.*}} = spv.SubgroupBallotKHR %{{.*}}: vector<4xi32>
     %0 = spv.SubgroupBallotKHR %predicate: vector<4xi32>
     spv.ReturnValue %0: vector<4xi32>
   }
 }

mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @iequal_scalar(%arg0: i32, %arg1: i32)  {
+  spv.func @iequal_scalar(%arg0: i32, %arg1: i32)  "None" {
     // CHECK: {{.*}} = spv.IEqual {{.*}}, {{.*}} : i32
     %0 = spv.IEqual %arg0, %arg1 : i32
     spv.Return
   }
-  func @inotequal_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) {
+  spv.func @inotequal_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.INotEqual {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.INotEqual %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @sgt_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) {
+  spv.func @sgt_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.SGreaterThan {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.SGreaterThan %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @sge_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) {
+  spv.func @sge_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.SGreaterThanEqual {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.SGreaterThanEqual %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @slt_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) {
+  spv.func @slt_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.SLessThan {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.SLessThan %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @slte_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) {
+  spv.func @slte_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.SLessThanEqual {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.SLessThanEqual %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @ugt_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) {
+  spv.func @ugt_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.UGreaterThan {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.UGreaterThan %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @ugte_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) {
+  spv.func @ugte_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.UGreaterThanEqual {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.UGreaterThanEqual %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @ult_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) {
+  spv.func @ult_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.ULessThan {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.ULessThan %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @ulte_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>)  {
+  spv.func @ulte_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>)  "None" {
     // CHECK: {{.*}} = spv.ULessThanEqual {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.ULessThanEqual %arg0, %arg1 : vector<4xi32>
     spv.Return
   }
-  func @cmpf(%arg0 : f32, %arg1 : f32) {
+  spv.func @cmpf(%arg0 : f32, %arg1 : f32) "None" {
     // CHECK: spv.FOrdEqual
     %1 = spv.FOrdEqual %arg0, %arg1 : f32
     // CHECK: spv.FOrdGreaterThan
     %2 = spv.FOrdGreaterThan %arg0, %arg1 : f32
     // CHECK: spv.FOrdGreaterThanEqual
     %3 = spv.FOrdGreaterThanEqual %arg0, %arg1 : f32
     // CHECK: spv.FOrdLessThan
     %4 = spv.FOrdLessThan %arg0, %arg1 : f32
spv.func @cmpf(%arg0 : f32, %arg1 : f32) "None" {
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.specConstant @condition_scalar = true
-  func @select() -> () {
+  spv.func @select() -> () "None" {
     %0 = spv.constant 4.0 : f32
     %1 = spv.constant 5.0 : f32
     %2 = spv._reference_of @condition_scalar : i1
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, f32
     %3 = spv.Select %2, %0, %1 : i1, f32
     %4 = spv.constant dense<[2.0, 3.0, 4.0, 5.0]> : vector<4xf32>
     %5 = spv.constant dense<[6.0, 7.0, 8.0, 9.0]> : vector<4xf32>
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, vector<4xf32>
     %6 = spv.Select %2, %4, %5 : i1, vector<4xf32>
     %7 = spv.constant dense<[true, true, true, true]> : vector<4xi1>
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : vector<4xi1>, vector<4xf32>
     %8 = spv.Select %7, %4, %5 : vector<4xi1>, vector<4xf32>
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/loop.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 // Single loop
 
 spv.module "Logical" "GLSL450" {
   // for (int i = 0; i < count; ++i) {}
-  func @loop(%count : i32) -> () {
+  spv.func @loop(%count : i32) -> () "None" {
     %zero = spv.constant 0: i32
     %one = spv.constant 1: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
 // CHECK:        spv.Branch ^bb1
 // CHECK-NEXT: ^bb1:
 // CHECK-NEXT:   spv.loop
     spv.loop {

 // CHECK-NEXT:   ^bb4:
 // CHECK-NEXT:     spv._merge
     ^merge:
       spv._merge
     }
     spv.Return
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 } attributes {
   capabilities = ["Shader"]
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @GV1 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
   spv.globalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
-  func @loop_kernel() {
+  spv.func @loop_kernel() "None" {
     %0 = spv._address_of @GV1 : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
     %1 = spv.constant 0 : i32
     %2 = spv.AccessChain %0[%1] : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
     %3 = spv._address_of @GV2 : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
     %5 = spv.AccessChain %3[%1] : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
     %6 = spv.constant 4 : i32
     %7 = spv.constant 42 : i32
     %8 = spv.constant 2 : i32

 // -----
 
 // Nested loop
 
 spv.module "Logical" "GLSL450" {
   // for (int i = 0; i < count; ++i) {
   //   for (int j = 0; j < count; ++j) { }
   // }
-  func @loop(%count : i32) -> () {
+  spv.func @loop(%count : i32) -> () "None" {
     %zero = spv.constant 0: i32
     %one = spv.constant 1: i32
     %ivar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     %jvar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
 // CHECK:        spv.Branch ^bb1
 // CHECK-NEXT: ^bb1:
 // CHECK-NEXT:   spv.loop

 // CHECK-NEXT:   ^bb5:
 // CHECK-NEXT:     spv._merge
     ^merge:
       spv._merge
     } // end outer loop
     spv.Return
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 } attributes {
   capabilities = ["Shader"]
 }
 

mlir/test/Dialect/SPIRV/Serialization/memory-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-// CHECK:           func {{@.*}}([[ARG1:%.*]]: !spv.ptr<f32, Input>, [[ARG2:%.*]]: !spv.ptr<f32, Output>) {
+// CHECK:           spv.func {{@.*}}([[ARG1:%.*]]: !spv.ptr<f32, Input>, [[ARG2:%.*]]: !spv.ptr<f32, Output>) "None" {
 // CHECK-NEXT:        [[VALUE:%.*]] = spv.Load "Input" [[ARG1]] : f32
 // CHECK-NEXT:        spv.Store "Output" [[ARG2]], [[VALUE]] : f32
 
 spv.module "Logical" "GLSL450" {
-  func @load_store(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) {
+  spv.func @load_store(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) "None" {
     %1 = spv.Load "Input" %arg0 : f32
     spv.Store "Output" %arg1, %1 : f32
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @access_chain(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>,
-                     %arg1 : i32, %arg2 : i32) {
+  spv.func @access_chain(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, %arg1 : i32, %arg2 : i32) "None" {
     // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, Function>
     // CHECK-NEXT: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, Function>
     %1 = spv.AccessChain %arg0[%arg1] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>
     %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/module.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 // CHECK:      spv.module "Logical" "GLSL450" {
-// CHECK-NEXT:   func @foo() {
+// CHECK-NEXT:   spv.func @foo() "None" {
 // CHECK-NEXT:     spv.Return
 // CHECK-NEXT:   }
 // CHECK-NEXT: } attributes {major_version = 1 : i32, minor_version = 0 : i32}
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
      spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
 } attributes {

mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
   // CHECK-LABEL: @group_non_uniform_ballot
-  func @group_non_uniform_ballot(%predicate: i1) -> vector<4xi32> {
+  spv.func @group_non_uniform_ballot(%predicate: i1) -> vector<4xi32> "None" {
     // CHECK: %{{.*}} = spv.GroupNonUniformBallot "Workgroup" %{{.*}}: vector<4xi32>
   %0 = spv.GroupNonUniformBallot "Workgroup" %predicate : vector<4xi32>
     spv.ReturnValue %0: vector<4xi32>
   }
 
   // CHECK-LABEL: @group_non_uniform_elect
-  func @group_non_uniform_elect() -> i1 {
+  spv.func @group_non_uniform_elect() -> i1 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformElect "Workgroup" : i1
     %0 = spv.GroupNonUniformElect "Workgroup" : i1
     spv.ReturnValue %0: i1
   }
 
   // CHECK-LABEL: @group_non_uniform_fadd_reduce
-  func @group_non_uniform_fadd_reduce(%val: f32) -> f32 {
+  spv.func @group_non_uniform_fadd_reduce(%val: f32) -> f32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformFAdd "Workgroup" "Reduce" %{{.+}} : f32
     %0 = spv.GroupNonUniformFAdd "Workgroup" "Reduce" %val : f32
     spv.ReturnValue %0: f32
   }
 
   // CHECK-LABEL: @group_non_uniform_fmax_reduce
-  func @group_non_uniform_fmax_reduce(%val: f32) -> f32 {
+  spv.func @group_non_uniform_fmax_reduce(%val: f32) -> f32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformFMax "Workgroup" "Reduce" %{{.+}} : f32
     %0 = spv.GroupNonUniformFMax "Workgroup" "Reduce" %val : f32
     spv.ReturnValue %0: f32
   }
 
   // CHECK-LABEL: @group_non_uniform_fmin_reduce
-  func @group_non_uniform_fmin_reduce(%val: f32) -> f32 {
+  spv.func @group_non_uniform_fmin_reduce(%val: f32) -> f32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformFMin "Workgroup" "Reduce" %{{.+}} : f32
     %0 = spv.GroupNonUniformFMin "Workgroup" "Reduce" %val : f32
     spv.ReturnValue %0: f32
   }
 
   // CHECK-LABEL: @group_non_uniform_fmul_reduce
-  func @group_non_uniform_fmul_reduce(%val: f32) -> f32 {
+  spv.func @group_non_uniform_fmul_reduce(%val: f32) -> f32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformFMul "Workgroup" "Reduce" %{{.+}} : f32
     %0 = spv.GroupNonUniformFMul "Workgroup" "Reduce" %val : f32
     spv.ReturnValue %0: f32
   }
 
   // CHECK-LABEL: @group_non_uniform_iadd_reduce
-  func @group_non_uniform_iadd_reduce(%val: i32) -> i32 {
+  spv.func @group_non_uniform_iadd_reduce(%val: i32) -> i32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformIAdd "Workgroup" "Reduce" %{{.+}} : i32
     %0 = spv.GroupNonUniformIAdd "Workgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 
   // CHECK-LABEL: @group_non_uniform_iadd_clustered_reduce
-  func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
+  spv.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> "None" {
     %four = spv.constant 4 : i32
     // CHECK: %{{.+}} = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>
     %0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32>
     spv.ReturnValue %0: vector<2xi32>
   }
 
   // CHECK-LABEL: @group_non_uniform_imul_reduce
-  func @group_non_uniform_imul_reduce(%val: i32) -> i32 {
+  spv.func @group_non_uniform_imul_reduce(%val: i32) -> i32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformIMul "Workgroup" "Reduce" %{{.+}} : i32
     %0 = spv.GroupNonUniformIMul "Workgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 
   // CHECK-LABEL: @group_non_uniform_smax_reduce
-  func @group_non_uniform_smax_reduce(%val: i32) -> i32 {
+  spv.func @group_non_uniform_smax_reduce(%val: i32) -> i32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformSMax "Workgroup" "Reduce" %{{.+}} : i32
     %0 = spv.GroupNonUniformSMax "Workgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 
   // CHECK-LABEL: @group_non_uniform_smin_reduce
-  func @group_non_uniform_smin_reduce(%val: i32) -> i32 {
+  spv.func @group_non_uniform_smin_reduce(%val: i32) -> i32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformSMin "Workgroup" "Reduce" %{{.+}} : i32
     %0 = spv.GroupNonUniformSMin "Workgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 
   // CHECK-LABEL: @group_non_uniform_umax_reduce
-  func @group_non_uniform_umax_reduce(%val: i32) -> i32 {
+  spv.func @group_non_uniform_umax_reduce(%val: i32) -> i32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformUMax "Workgroup" "Reduce" %{{.+}} : i32
     %0 = spv.GroupNonUniformUMax "Workgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 
   // CHECK-LABEL: @group_non_uniform_umin_reduce
-  func @group_non_uniform_umin_reduce(%val: i32) -> i32 {
+  spv.func @group_non_uniform_umin_reduce(%val: i32) -> i32 "None" {
     // CHECK: %{{.+}} = spv.GroupNonUniformUMin "Workgroup" "Reduce" %{{.+}} : i32
     %0 = spv.GroupNonUniformUMin "Workgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 }

mlir/test/Dialect/SPIRV/Serialization/phi.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 // Test branch with one block argument
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
 // CHECK:        %[[CST:.*]] = spv.constant 0
     %zero = spv.constant 0 : i32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[CST]] : i32)
     spv.Branch ^bb1(%zero : i32)
 // CHECK-NEXT: ^bb1(%{{.*}}: i32):
   ^bb1(%arg0: i32):
    spv.Return
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 } attributes {
   capabilities = ["Shader"]
 }
 
 // -----
 
 // Test branch with multiple block arguments
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
 // CHECK:        %[[ZERO:.*]] = spv.constant 0
     %zero = spv.constant 0 : i32
 // CHECK-NEXT:   %[[ONE:.*]] = spv.constant 1
     %one = spv.constant 1.0 : f32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[ZERO]], %[[ONE]] : i32, f32)
     spv.Branch ^bb1(%zero, %one : i32, f32)
 
 // CHECK-NEXT: ^bb1(%{{.*}}: i32, %{{.*}}: f32):     // pred: ^bb0
   ^bb1(%arg0: i32, %arg1: f32):
    spv.Return
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 } attributes {
   capabilities = ["Shader"]
 }
 
 // -----
 
 // Test using block arguments within branch
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
 // CHECK:        %[[CST0:.*]] = spv.constant 0
     %zero = spv.constant 0 : i32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[CST0]] : i32)
     spv.Branch ^bb1(%zero : i32)
 
 // CHECK-NEXT: ^bb1(%[[ARG:.*]]: i32):
   ^bb1(%arg0: i32):
 // CHECK-NEXT:   %[[ADD:.*]] = spv.IAdd %[[ARG]], %[[ARG]] : i32
     %0 = spv.IAdd %arg0, %arg0 : i32
 // CHECK-NEXT:   %[[CST1:.*]] = spv.constant 0
 // CHECK-NEXT:   spv.Branch ^bb2(%[[CST1]], %[[ADD]] : i32, i32)
     spv.Branch ^bb2(%zero, %0 : i32, i32)
 
 // CHECK-NEXT: ^bb2(%{{.*}}: i32, %{{.*}}: i32):
   ^bb2(%arg1: i32, %arg2: i32):
    spv.Return
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 } attributes {
   capabilities = ["Shader"]
 }
 
 // -----
 
 // Test block not following domination order
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
 // CHECK:        spv.Branch ^bb1
     spv.Branch ^bb1
 
 // CHECK-NEXT: ^bb1:
 // CHECK-NEXT:   %[[ZERO:.*]] = spv.constant 0
 // CHECK-NEXT:   %[[ONE:.*]] = spv.constant 1
 // CHECK-NEXT:   spv.Branch ^bb2(%[[ZERO]], %[[ONE]] : i32, f32)
 
 // CHECK-NEXT: ^bb2(%{{.*}}: i32, %{{.*}}: f32):
   ^bb2(%arg0: i32, %arg1: f32):
 // CHECK-NEXT:   spv.Return
    spv.Return
 
   // This block is reordered to follow domination order.
   ^bb1:
     %zero = spv.constant 0 : i32
     %one = spv.constant 1.0 : f32
     spv.Branch ^bb2(%zero, %one : i32, f32)
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 } attributes {
   capabilities = ["Shader"]
 }
 
 // -----
 
 // Test multiple predecessors
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
     %var = spv.Variable : !spv.ptr<i32, Function>
 
 // CHECK:      spv.selection
     spv.selection {
       %true = spv.constant true
 // CHECK:        spv.BranchConditional %{{.*}}, ^bb1, ^bb2
       spv.BranchConditional %true, ^true, ^false
 

 // CHECK-NEXT: ^bb4:
     ^merge:
 // CHECK-NEXT:   spv._merge
       spv._merge
     }
     spv.Return
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 } attributes {
   capabilities = ["Shader"]
 }
 
 // -----
 
 // Test nested loops with block arguments
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
   spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
-  func @fmul_kernel() {
+  spv.func @fmul_kernel() "None" {
     %3 = spv.constant 12 : i32
     %4 = spv.constant 32 : i32
     %5 = spv.constant 4 : i32
     %6 = spv._address_of @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %7 = spv.Load "Input" %6 : vector<3xi32>
     %8 = spv.CompositeExtract %7[0 : i32] : vector<3xi32>
     %9 = spv._address_of @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %10 = spv.Load "Input" %9 : vector<3xi32>

mlir/test/Dialect/SPIRV/Serialization/selection.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 // Selection with both then and else branches
 
 spv.module "Logical" "GLSL450" {
-  func @selection(%cond: i1) -> () {
+  spv.func @selection(%cond: i1) -> () "None" {
 // CHECK:        spv.Branch ^bb1
 // CHECK-NEXT: ^bb1:
     %zero = spv.constant 0: i32
     %one = spv.constant 1: i32
     %two = spv.constant 2: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
 // CHECK-NEXT:   spv.selection {
// CHECK-NEXT:   ^bb3:
     ^merge:
 // CHECK-NEXT:     spv._merge
       spv._merge
     }
 
     spv.Return
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
   spv.ExecutionMode @main "LocalSize", 1, 1, 1
 } attributes {
   capabilities = ["Shader"]
 }
 
 // -----
 
 // Selection with only then branch
 // Selection in function entry block
 
 spv.module "Logical" "GLSL450" {
-// CHECK:      func @selection(%[[ARG:.*]]: i1
-  func @selection(%cond: i1) -> (i32) {
+// CHECK:      spv.func @selection(%[[ARG:.*]]: i1
+  spv.func @selection(%cond: i1) -> (i32) "None" {
 // CHECK:        spv.Branch ^bb1
 // CHECK-NEXT: ^bb1:
 // CHECK-NEXT:   spv.selection
     spv.selection {
 // CHECK-NEXT: spv.BranchConditional %[[ARG]], ^bb1, ^bb2
       spv.BranchConditional %cond, ^then, ^merge
 
 // CHECK:        ^bb1:
     ^then:
       %zero = spv.constant 0 : i32
       spv.ReturnValue  %zero : i32
 
 // CHECK:        ^bb2:
     ^merge:
 // CHECK-NEXT:     spv._merge
       spv._merge
     }
 
     %one = spv.constant 1 : i32
     spv.ReturnValue  %one : i32
   }
 
-  func @main() -> () {
+  spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
   spv.ExecutionMode @main "LocalSize", 1, 1, 1
 } attributes {
   capabilities = ["Shader"]
 }
 

mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
   // CHECK: spv.specConstant @sc_true = true
   spv.specConstant @sc_true = true
   // CHECK: spv.specConstant @sc_false spec_id(1) = false
   spv.specConstant @sc_false spec_id(1) = false
 
   // CHECK: spv.specConstant @sc_int = -5 : i32
   spv.specConstant @sc_int = -5 : i32
 
   // CHECK: spv.specConstant @sc_float spec_id(5) = 1.000000e+00 : f32
   spv.specConstant @sc_float spec_id(5) = 1. : f32
 
   // CHECK-LABEL: @use
-  func @use() -> (i32) {
+  spv.func @use() -> (i32) "None" {
     // We materialize a `spv._reference_of` op at every use of a
     // specialization constant in the deserializer. So two ops here.
     // CHECK: %[[USE1:.*]] = spv._reference_of @sc_int : i32
     // CHECK: %[[USE2:.*]] = spv._reference_of @sc_int : i32
     // CHECK: spv.IAdd %[[USE1]], %[[USE2]]
 
     %0 = spv._reference_of @sc_int : i32
     %1 = spv.IAdd %0, %0 : i32
     spv.ReturnValue %1 : i32
   }
 }

mlir/test/Dialect/SPIRV/Serialization/struct.mlir

spv.module "Logical" "GLSL450" {
   // CHECK: !spv.ptr<!spv.struct<f32 [0, NonWritable], i32 [4, NonWritable, NonReadable]>, StorageBuffer>
   spv.globalVariable @var6 : !spv.ptr<!spv.struct<f32 [0, NonWritable], i32 [4, NonWritable, NonReadable]>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<>, StorageBuffer>
   spv.globalVariable @empty : !spv.ptr<!spv.struct<>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>,
   // CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Output>
-  func @kernel_1(%arg0: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>, %arg1: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Output>) -> () {
+  spv.func @kernel_1(%arg0: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>, %arg1: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Output>) -> () "None" {
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/terminator.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
   // CHECK-LABEL: @ret
-  func @ret() -> () {
+  spv.func @ret() -> () "None" {
     // CHECK: spv.Return
     spv.Return
   }
 
   // CHECK-LABEL: @ret_val
-  func @ret_val() -> (i32) {
+  spv.func @ret_val() -> (i32) "None" {
     %0 = spv.Variable : !spv.ptr<i32, Function>
     %1 = spv.Load "Function" %0 : i32
     // CHECK: spv.ReturnValue {{.*}} : i32
     spv.ReturnValue %1 : i32
   }
 
   // CHECK-LABEL: @unreachable
-  func @unreachable() {
+  spv.func @unreachable() "None" {
     spv.Return
   // CHECK-NOT: ^bb
   ^bb1:
     // Unreachable blocks will be dropped during serialization.
     // CHECK-NOT: spv.Unreachable
     spv.Unreachable
   }
 }

mlir/test/Dialect/SPIRV/Serialization/undef.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
     // CHECK: {{%.*}} = spv.undef : f32
     // CHECK-NEXT: {{%.*}} = spv.undef : f32
     %0 = spv.undef : f32
     %1 = spv.undef : f32
     %2 = spv.FAdd %0, %1 : f32
     // CHECK: {{%.*}} = spv.undef : vector<4xi32>
     %3 = spv.undef : vector<4xi32>
     %4 = spv.CompositeExtract %3[1 : i32] : vector<4xi32>
     // CHECK: {{%.*}} = spv.undef : !spv.array<4 x !spv.array<4 x i32>>
     %5 = spv.undef : !spv.array<4x!spv.array<4xi32>>
     %6 = spv.CompositeExtract %5[1 : i32, 2 : i32] : !spv.array<4x!spv.array<4xi32>>
     // CHECK: {{%.*}} = spv.undef : !spv.ptr<!spv.struct<f32>, StorageBuffer>
     %7 = spv.undef : !spv.ptr<!spv.struct<f32>, StorageBuffer>
     %8 = spv.constant 0 : i32
     %9 = spv.AccessChain %7[%8] : !spv.ptr<!spv.struct<f32>, StorageBuffer>
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  // CHECK: func {{@.*}}
-  func @ignore_unused_undef() -> () {
+  // CHECK: spv.func {{@.*}}
+  spv.func @ignore_unused_undef() -> () "None" {
     // CHECK-NEXT: spv.Return
     %0 = spv.undef : f32
     spv.Return
   }
-}
+}
 

mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir

spv.module "Logical" "GLSL450" {
   spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
   // CHECK-DAG: spv.globalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR2:@.*]] bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR3:@.*]] bind(0, 3) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR4:@.*]] bind(0, 4) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR5:@.*]] bind(0, 5) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR6:@.*]] bind(0, 6) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
-  // CHECK: func [[FN:@.*]]()
-  func @load_store_kernel(%arg0: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32>>>, StorageBuffer>
-                          {spv.interface_var_abi = {binding = 0 : i32,
-                                                    descriptor_set = 0 : i32,
-                                                    storage_class = 12 : i32}},
-                          %arg1: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32>>>, StorageBuffer>
-                          {spv.interface_var_abi = {binding = 1 : i32,
-                                                    descriptor_set = 0 : i32,
-                                                    storage_class = 12 : i32}},
-                          %arg2: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32>>>, StorageBuffer>
-                          {spv.interface_var_abi = {binding = 2 : i32,
-                                                    descriptor_set = 0 : i32,
-                                                    storage_class = 12 : i32}},
-                          %arg3: i32
-                          {spv.interface_var_abi = {binding = 3 : i32,
-                                                    descriptor_set = 0 : i32,
-                                                    storage_class = 12 : i32}},
-                          %arg4: i32
-                          {spv.interface_var_abi = {binding = 4 : i32,
-                                                    descriptor_set = 0 : i32,
-                                                    storage_class = 12 : i32}},
-                          %arg5: i32
-                          {spv.interface_var_abi = {binding = 5 : i32,
-                                                    descriptor_set = 0 : i32,
-                                                    storage_class = 12 : i32}},
-                          %arg6: i32
-                          {spv.interface_var_abi = {binding = 6 : i32,
-                                                    descriptor_set = 0 : i32,
-                                                    storage_class = 12 : i32}})
+  // CHECK: spv.func [[FN:@.*]]()
+  spv.func @load_store_kernel(
+    %arg0: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32>>>, StorageBuffer>
+                              storage_class = 12 : i32}}) "None"
   attributes  {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     // CHECK: [[ADDRESSARG6:%.*]] = spv._address_of [[VAR6]]
     // CHECK: [[CONST6:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]
     // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG6PTR]]
     // CHECK: [[ADDRESSARG5:%.*]] = spv._address_of [[VAR5]]
     // CHECK: [[CONST5:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG5PTR:%.*]] = spv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]]

mlir/test/Dialect/SPIRV/Transforms/abi-simple.mlir

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 // CHECK-LABEL: spv.module
 spv.module "Logical" "GLSL450" {
   // CHECK-DAG:    spv.globalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<f32 [0]>, StorageBuffer>
   // CHECK-DAG:    spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
-  // CHECK:    func [[FN:@.*]]()
-  func @kernel_1(%arg0: f32
+  // CHECK:    spv.func [[FN:@.*]]()
+  spv.func @kernel(%arg0: f32
                 {spv.interface_var_abi = {binding = 0 : i32,
                                           descriptor_set = 0 : i32,
                                           storage_class = 12 : i32}},
                  %arg1: !spv.ptr<!spv.struct<!spv.array<12 x f32>>, StorageBuffer>
                  {spv.interface_var_abi = {binding = 1 : i32,
                                            descriptor_set = 0 : i32,
-                                           storage_class = 12 : i32}})
-  attributes  {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
+                                           storage_class = 12 : i32}}) "None"
+  attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     // CHECK: [[ARG1:%.*]] = spv._address_of [[VAR1]]
     // CHECK: [[ADDRESSARG0:%.*]] = spv._address_of [[VAR0]]
     // CHECK: [[CONST0:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG0PTR:%.*]] = spv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]]
     // CHECK: [[ARG0:%.*]] = spv.Load "StorageBuffer" [[ARG0PTR]]
     // CHECK: spv.Return
     spv.Return
   }
   // CHECK: spv.EntryPoint "GLCompute" [[FN]]
   // CHECK: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
 } attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}

mlir/test/Dialect/SPIRV/Transforms/inlining.mlir

 // RUN: mlir-opt %s -split-input-file -pass-pipeline='spv.module(inline)' -mlir-disable-inline-simplify | FileCheck %s
 
 spv.module "Logical" "GLSL450" {
-  func @callee() {
+  spv.func @callee() "None" {
     spv.Return
   }
 
-  // CHECK-LABEL: func @calling_single_block_ret_func
-  func @calling_single_block_ret_func() {
+  // CHECK-LABEL: @calling_single_block_ret_func
+  spv.func @calling_single_block_ret_func() "None" {
     // CHECK-NEXT: spv.Return
     spv.FunctionCall @callee() : () -> ()
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @callee() -> i32 {
+  spv.func @callee() -> i32 "None" {
     %0 = spv.constant 42 : i32
     spv.ReturnValue %0 : i32
   }
 
-  // CHECK-LABEL: func @calling_single_block_retval_func
-  func @calling_single_block_retval_func() -> i32 {
+  // CHECK-LABEL: @calling_single_block_retval_func
+  spv.func @calling_single_block_retval_func() -> i32 "None" {
     // CHECK-NEXT: %[[CST:.*]] = spv.constant 42
     %0 = spv.FunctionCall @callee() : () -> (i32)
     // CHECK-NEXT: spv.ReturnValue %[[CST]]
     spv.ReturnValue %0 : i32
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @data bind(0, 0) : !spv.ptr<!spv.struct<!spv.rtarray<i32> [0]>, StorageBuffer>
-  func @callee() {
+  spv.func @callee() "None" {
     %0 = spv._address_of @data : !spv.ptr<!spv.struct<!spv.rtarray<i32> [0]>, StorageBuffer>
     %1 = spv.constant 0: i32
     %2 = spv.AccessChain %0[%1, %1] : !spv.ptr<!spv.struct<!spv.rtarray<i32> [0]>, StorageBuffer>
     spv.Branch ^next
 
   ^next:
     %3 = spv.constant 42: i32
     spv.Store "StorageBuffer" %2, %3 : i32
     spv.Return
   }
 
-  // CHECK-LABEL: func @calling_multi_block_ret_func
-  func @calling_multi_block_ret_func() {
+  // CHECK-LABEL: @calling_multi_block_ret_func
+  spv.func @calling_multi_block_ret_func() "None" {
     // CHECK-NEXT:   spv._address_of
     // CHECK-NEXT:   spv.constant 0
     // CHECK-NEXT:   spv.AccessChain
     // CHECK-NEXT:   spv.Branch ^bb1
     // CHECK-NEXT: ^bb1:
     // CHECK-NEXT:   spv.constant
     // CHECK-NEXT:   spv.Store
     // CHECK-NEXT:   spv.Branch ^bb2
     spv.FunctionCall @callee() : () -> ()
     // CHECK-NEXT: ^bb2:
     // CHECK-NEXT:   spv.Return
     spv.Return
   }
 }
 
 // TODO: calling_multi_block_retval_func
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @callee(%cond : i1) -> () {
+  spv.func @callee(%cond : i1) -> () "None" {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
       spv.Return
     ^merge:
       spv._merge
     }
     spv.Return
   }
 
-  // CHECK-LABEL: calling_selection_ret_func
-  func @calling_selection_ret_func() {
+  // CHECK-LABEL: @calling_selection_ret_func
+  spv.func @calling_selection_ret_func() "None" {
     %0 = spv.constant true
     // CHECK: spv.FunctionCall
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @callee(%cond : i1) -> () {
+  spv.func @callee(%cond : i1) -> () "None" {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
       spv.Branch ^merge
     ^merge:
       spv._merge
     }
     spv.Return
   }
 
-  // CHECK-LABEL: calling_selection_no_ret_func
-  func @calling_selection_no_ret_func() {
+  // CHECK-LABEL: @calling_selection_no_ret_func
+  spv.func @calling_selection_no_ret_func() "None" {
     // CHECK-NEXT: %[[TRUE:.*]] = spv.constant true
     %0 = spv.constant true
     // CHECK-NEXT: spv.selection
     // CHECK-NEXT:   spv.BranchConditional %[[TRUE]], ^bb1, ^bb2
     // CHECK-NEXT: ^bb1:
     // CHECK-NEXT:   spv.Branch ^bb2
     // CHECK-NEXT: ^bb2:
     // CHECK-NEXT:   spv._merge
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @callee(%cond : i1) -> () {
+  spv.func @callee(%cond : i1) -> () "None" {
     spv.loop {
       spv.Branch ^header
     ^header:
       spv.BranchConditional %cond, ^body, ^merge
     ^body:
       spv.Return
     ^continue:
       spv.Branch ^header
     ^merge:
       spv._merge
     }
     spv.Return
   }
 
-  // CHECK-LABEL: calling_loop_ret_func
-  func @calling_loop_ret_func() {
+  // CHECK-LABEL: @calling_loop_ret_func
+  spv.func @calling_loop_ret_func() "None" {
     %0 = spv.constant true
     // CHECK: spv.FunctionCall
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @callee(%cond : i1) -> () {
+  spv.func @callee(%cond : i1) -> () "None" {
     spv.loop {
       spv.Branch ^header
     ^header:
       spv.BranchConditional %cond, ^body, ^merge
     ^body:
       spv.Branch ^continue
     ^continue:
       spv.Branch ^header
     ^merge:
       spv._merge
     }
     spv.Return
   }
 
-  // CHECK-LABEL: calling_loop_no_ret_func
-  func @calling_loop_no_ret_func() {
+  // CHECK-LABEL: @calling_loop_no_ret_func
+  spv.func @calling_loop_no_ret_func() "None" {
     // CHECK-NEXT: %[[TRUE:.*]] = spv.constant true
     %0 = spv.constant true
     // CHECK-NEXT: spv.loop
     // CHECK-NEXT:   spv.Branch ^bb1
     // CHECK-NEXT: ^bb1:
     // CHECK-NEXT:   spv.BranchConditional %[[TRUE]], ^bb2, ^bb4
     // CHECK-NEXT: ^bb2:
     // CHECK-NEXT:   spv.Branch ^bb3
     // CHECK-NEXT: ^bb3:
     // CHECK-NEXT:   spv.Branch ^bb1
     // CHECK-NEXT: ^bb4:
     // CHECK-NEXT:   spv._merge
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   spv.globalVariable @arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
-  // CHECK: func @inline_into_selection_region
-  func @inline_into_selection_region() {
+
+  // CHECK: @inline_into_selection_region
+  spv.func @inline_into_selection_region() "None" {
     %1 = spv.constant 0 : i32
     // CHECK-DAG: [[ADDRESS_ARG0:%.*]] = spv._address_of @arg_0
     // CHECK-DAG: [[ADDRESS_ARG1:%.*]] = spv._address_of @arg_1
     // CHECK-DAG: [[LOADPTR:%.*]] = spv.AccessChain [[ADDRESS_ARG0]]
     // CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOADPTR]]
     %2 = spv._address_of @arg_0 : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
     %3 = spv._address_of @arg_1 : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
     %4 = spv.AccessChain %2[%1] : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
^bb1: // pred: ^bb0
       spv.FunctionCall @atomic_add(%5, %7) : (i32, !spv.ptr<i32, StorageBuffer>) -> ()
       spv.Branch ^bb2
     ^bb2 : // 2 preds: ^bb0, ^bb1
       spv._merge
     }
     // CHECK: spv.Return
     spv.Return
   }
-  func @atomic_add(%arg0: i32, %arg1: !spv.ptr<i32, StorageBuffer>) {
+  spv.func @atomic_add(%arg0: i32, %arg1: !spv.ptr<i32, StorageBuffer>) "None" {
     %0 = spv.AtomicIAdd "Device" "AcquireRelease" %arg1, %arg0 : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }
   spv.EntryPoint "GLCompute" @inline_into_selection_region
   spv.ExecutionMode @inline_into_selection_region "LocalSize", 32, 1, 1
 } attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
 
 // TODO: Add tests for inlining structured control flow into
-// structured control flow.
+// structured control flow.
 

mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir

spv.module "Logical" "GLSL450" {
   spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.array<16x!spv.struct<f32, f32, !spv.array<16xf32>>>, f32>, StorageBuffer>
 
   // CHECK: spv.globalVariable @var4 bind(1, 2) : !spv.ptr<!spv.struct<!spv.struct<!spv.struct<i1 [0], i8 [1], i16 [2], i32 [4], i64 [8]> [0], f32 [16], i1 [20]> [0], i1 [24]>, StorageBuffer>
   spv.globalVariable @var4 bind(1,2) : !spv.ptr<!spv.struct<!spv.struct<!spv.struct<i1, i8, i16, i32, i64>, f32, i1>, i1>, StorageBuffer>
 
   // CHECK: spv.globalVariable @var5 bind(1, 3) : !spv.ptr<!spv.struct<!spv.array<256 x f32 [4]> [0]>, StorageBuffer>
   spv.globalVariable @var5 bind(1,3) : !spv.ptr<!spv.struct<!spv.array<256xf32>>, StorageBuffer>
 
-  func @kernel() -> () {
+  spv.func @kernel() -> () "None" {
     %c0 = spv.constant 0 : i32
     // CHECK: {{%.*}} = spv._address_of @var0 : !spv.ptr<!spv.struct<i32 [0], !spv.struct<f32 [0], i32 [4]> [4], f32 [12]>, Uniform>
     %0 = spv._address_of @var0 : !spv.ptr<!spv.struct<i32, !spv.struct<f32, i32>, f32>, Uniform>
     // CHECK:  {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.struct<i32 [0], !spv.struct<f32 [0], i32 [4]> [4], f32 [12]>, Uniform>
     %1 = spv.AccessChain %0[%c0] : !spv.ptr<!spv.struct<i32, !spv.struct<f32, i32>, f32>, Uniform>
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/control-flow-ops.mlir

 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.FunctionCall
 //===----------------------------------------------------------------------===//
 
 spv.module "Logical" "GLSL450" {
-  func @fmain(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>, %arg2 : i32) -> i32 {
+  spv.func @fmain(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>, %arg2 : i32) -> i32 "None" {
     // CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}, {{%.*}}) : (vector<4xf32>, vector<4xf32>) -> vector<4xf32>
     %0 = spv.FunctionCall @f_0(%arg0, %arg1) : (vector<4xf32>, vector<4xf32>) -> vector<4xf32>
     // CHECK: spv.FunctionCall @f_1({{%.*}}, {{%.*}}) : (vector<4xf32>, vector<4xf32>) -> ()
     spv.FunctionCall @f_1(%0, %arg1) : (vector<4xf32>, vector<4xf32>) ->  ()
     // CHECK: spv.FunctionCall @f_2() : () -> ()
     spv.FunctionCall @f_2() : () -> ()
     // CHECK: {{%.*}} = spv.FunctionCall @f_3({{%.*}}) : (i32) -> i32
     %1 = spv.FunctionCall @f_3(%arg2) : (i32) -> i32
     spv.ReturnValue %1 : i32
   }
 
-  func @f_0(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) -> (vector<4xf32>) {
+  spv.func @f_0(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) -> (vector<4xf32>) "None" {
     spv.ReturnValue %arg0 : vector<4xf32>
   }
 
-  func @f_1(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) -> () {
+  spv.func @f_1(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) -> () "None" {
     spv.Return
   }
 
-  func @f_2() -> () {
+  spv.func @f_2() -> () "None" {
     spv.Return
   }
 
-  func @f_3(%arg0 : i32) -> (i32) {
+  spv.func @f_3(%arg0 : i32) -> (i32) "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
 // -----
 
 // Allow calling functions in other module-like ops
-func @callee() {
+spv.func @callee() "None" {
   spv.Return
 }
 
 func @caller() {
   // CHECK: spv.FunctionCall
   spv.FunctionCall @callee() : () -> ()
   spv.Return
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @f_invalid_result_type(%arg0 : i32, %arg1 : i32) -> () {
+  spv.func @f_invalid_result_type(%arg0 : i32, %arg1 : i32) -> () "None" {
     // expected-error @+1 {{expected callee function to have 0 or 1 result, but provided 2}}
     %0:2 = spv.FunctionCall @f_invalid_result_type(%arg0, %arg1) : (i32, i32) -> (i32, i32)
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @f_result_type_mismatch(%arg0 : i32, %arg1 : i32) -> () {
+  spv.func @f_result_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" {
     // expected-error @+1 {{has incorrect number of results has for callee: expected 0, but provided 1}}
     %1 = spv.FunctionCall @f_result_type_mismatch(%arg0, %arg0) : (i32, i32) -> (i32)
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () {
+  spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" {
     // expected-error @+1 {{has incorrect number of operands for callee: expected 2, but provided 1}}
     spv.FunctionCall @f_type_mismatch(%arg0) : (i32) -> ()
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () {
+  spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" {
     %0 = spv.constant 2.0 : f32
     // expected-error @+1 {{operand type mismatch: expected operand type 'i32', but provided 'f32' for operand number 1}}
     spv.FunctionCall @f_type_mismatch(%arg0, %0) : (i32, f32) -> ()
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> i32 {
+  spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> i32 "None" {
+    %cst = spv.constant 0: i32
     // expected-error @+1 {{result type mismatch: expected 'i32', but provided 'f32'}}
     %0 = spv.FunctionCall @f_type_mismatch(%arg0, %arg0) : (i32, i32) -> f32
-    spv.Return
+    spv.ReturnValue %cst: i32
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @f_foo(%arg0 : i32, %arg1 : i32) -> i32 {
+  spv.func @f_foo(%arg0 : i32, %arg1 : i32) -> i32 "None" {
     // expected-error @+1 {{op callee function 'f_undefined' not found in nearest symbol table}}
     %0 = spv.FunctionCall @f_undefined(%arg0, %arg0) : (i32, i32) -> i32
-    spv.Return
+    spv.ReturnValue %0: i32
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.loop
 //===----------------------------------------------------------------------===//
func @in_loop(%cond : i1) -> () {
   ^continue:
     spv.Branch ^header
   ^merge:
     spv._merge
   }
   spv.Return
 }
 
-// -----
-
 // CHECK-LABEL: in_other_func_like_op
 func @in_other_func_like_op() {
-  // CHECK: spv.Return 
+  // CHECK: spv.Return
   spv.Return
 }
 
 // -----
 
 "foo.function"() ({
   // expected-error @+1 {{op must appear in a function-like op's block}}
   spv.Return
 })  : () -> ()
 
 // -----
 
 // Return mismatches function signature
 spv.module "Logical" "GLSL450" {
-  func @work() -> (i32) {
+  spv.func @work() -> (i32) "None" {
     // expected-error @+1 {{cannot be used in functions returning value}}
     spv.Return
   }
 }
 
 // -----
 
+spv.module "Logical" "GLSL450" {
+  spv.func @in_nested_region(%cond: i1) -> (i32) "None" {
+    spv.selection {
+      spv.BranchConditional %cond, ^then, ^merge
+    ^then:
+      // expected-error @+1 {{cannot be used in functions returning value}}
+      spv.Return
+    ^merge:
+      spv._merge
+    }
+
+    %zero = spv.constant 0: i32
+    spv.ReturnValue %zero: i32
+  }
+}
+
+// -----
+
 //===----------------------------------------------------------------------===//
 // spv.ReturnValue
 //===----------------------------------------------------------------------===//
 
 func @ret_val() -> (i32) {
   %0 = spv.constant 42 : i32
   // CHECK: spv.ReturnValue %{{.*}} : i32
   spv.ReturnValue %0 : i32
^continue:
     spv.Branch ^header
   ^merge:
     spv._merge
   }
   %one = spv.constant 1 : i32
   spv.ReturnValue %one : i32
 }
 
+// CHECK-LABEL: in_other_func_like_op
+func @in_other_func_like_op(%arg: i32) -> i32 {
+  // CHECK: spv.ReturnValue
+  spv.ReturnValue %arg: i32
+}
+
 // -----
 
 "foo.function"() ({
   %0 = spv.constant true
   // expected-error @+1 {{op must appear in a function-like op's block}}
   spv.ReturnValue %0 : i1
 })  : () -> ()
 
 // -----
 
-func @value_count_mismatch() -> () {
+spv.module "Logical" "GLSL450" {
+  spv.func @value_count_mismatch() -> () "None" {
     %0 = spv.constant 42 : i32
     // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}}
     spv.ReturnValue %0 : i32
   }
+}
 
 // -----
 
-func @value_type_mismatch() -> (f32) {
+spv.module "Logical" "GLSL450" {
+  spv.func @value_type_mismatch() -> (f32) "None" {
     %0 = spv.constant 42 : i32
     // expected-error @+1 {{return value's type ('i32') mismatch with function's result type ('f32')}}
     spv.ReturnValue %0 : i32
   }
+}
+
+// -----
+
+spv.module "Logical" "GLSL450" {
+  spv.func @in_nested_region(%cond: i1) -> () "None" {
+    spv.selection {
+      spv.BranchConditional %cond, ^then, ^merge
+    ^then:
+      %cst = spv.constant 0: i32
+      // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}}
+      spv.ReturnValue %cst: i32
+    ^merge:
+      spv._merge
+    }
+
+    spv.Return
+  }
+}
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.selection
 //===----------------------------------------------------------------------===//
 
 func @selection(%cond: i1) -> () {

mlir/test/Dialect/SPIRV/ops.mlir

 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.ExecutionMode
 //===----------------------------------------------------------------------===//
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // CHECK: spv.ExecutionMode {{@.*}} "ContractionOff"
    spv.ExecutionMode @do_nothing "ContractionOff"
 }
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // CHECK: spv.ExecutionMode {{@.*}} "LocalSizeHint", 3, 4, 5
    spv.ExecutionMode @do_nothing "LocalSizeHint", 3, 4, 5
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // expected-error @+1 {{custom op 'spv.ExecutionMode' invalid execution_mode attribute specification: "GLCompute"}}
    spv.ExecutionMode @do_nothing "GLCompute", 3, 4, 5
 }
 
 // -----
func @aligned_load_incorrect_attributes() -> () {
   return
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
   // CHECK_LABEL: @simple_load
-  func @simple_load() -> () {
+  spv.func @simple_load() -> () "None" {
     // CHECK: spv.Load "Input" {{%.*}} : f32
     %0 = spv._address_of @var0 : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.Return
   }
 }
 
 // -----
func @aligned_store_incorrect_attributes(%arg0 : f32) -> () {
   spv.Store  "Function" %0, %arg0 ["Aligned", 4, 23] : f32
   return
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
-  func @simple_store(%arg0 : f32) -> () {
+  spv.func @simple_store(%arg0 : f32) -> () "None" {
     %0 = spv._address_of @var0 : !spv.ptr<f32, Input>
     // CHECK: spv.Store  "Input" {{%.*}}, {{%.*}} : f32
     spv.Store  "Input" %0, %arg0 : f32
     spv.Return
   }
 }
 
 // -----
func @variable_init_normal_constant() -> () {
   %1 = spv.Variable init(%0) : !spv.ptr<f32, Function>
   return
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @global : !spv.ptr<f32, Workgroup>
-  func @variable_init_global_variable() -> () {
+  spv.func @variable_init_global_variable() -> () "None" {
     %0 = spv._address_of @global : !spv.ptr<f32, Workgroup>
     // CHECK: spv.Variable init({{.*}}) : !spv.ptr<!spv.ptr<f32, Workgroup>, Function>
     %1 = spv.Variable init(%0) : !spv.ptr<!spv.ptr<f32, Workgroup>, Function>
     spv.Return
   }
 } attributes {
   capability = ["VariablePointers"],
   extension = ["SPV_KHR_variable_pointers"]
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.specConstant @sc = 42 : i32
   // CHECK-LABEL: @variable_init_spec_constant
-  func @variable_init_spec_constant() -> () {
+  spv.func @variable_init_spec_constant() -> () "None" {
     %0 = spv._reference_of @sc : i32
     // CHECK: spv.Variable init(%0) : !spv.ptr<i32, Function>
     %1 = spv.Variable init(%0) : !spv.ptr<i32, Function>
     spv.Return
   }
 }
 
 // -----

mlir/test/Dialect/SPIRV/structure-ops.mlir

 // RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spv._address_of
 //===----------------------------------------------------------------------===//
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
-  func @access_chain() -> () {
+  spv.func @access_chain() -> () "None" {
     %0 = spv.constant 1: i32
     // CHECK: [[VAR1:%.*]] = spv._address_of @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4 x f32>>, Input>
     // CHECK-NEXT: spv.AccessChain [[VAR1]][{{.*}}, {{.*}}] : !spv.ptr<!spv.struct<f32, !spv.array<4 x f32>>, Input>
     %1 = spv._address_of @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
     %2 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
     spv.Return
   }
 }
 
 // -----
 
 // Allow taking address of global variables in other module-like ops
 spv.globalVariable @var : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
 func @address_of() -> () {
   // CHECK: spv._address_of @var
   %1 = spv._address_of @var : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
   return
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
     // expected-error @+1 {{expected spv.globalVariable symbol}}
     %0 = spv._address_of @var2 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced global variable's type}}
     %0 = spv._address_of @var1 : !spv.ptr<f32, Input>
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//

 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.EntryPoint
 //===----------------------------------------------------------------------===//
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // CHECK: spv.EntryPoint "GLCompute" @do_nothing
    spv.EntryPoint "GLCompute" @do_nothing
 }
 
 spv.module "Logical" "GLSL450" {
    spv.globalVariable @var2 : !spv.ptr<f32, Input>
    spv.globalVariable @var3 : !spv.ptr<f32, Output>
-   func @do_something(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () {
+   spv.func @do_something(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () "None" {
      %1 = spv.Load "Input" %arg0 : f32
      spv.Store "Output" %arg1, %1 : f32
      spv.Return
    }
    // CHECK: spv.EntryPoint "GLCompute" @do_something, @var2, @var3
    spv.EntryPoint "GLCompute" @do_something, @var2, @var3
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // expected-error @+1 {{invalid kind of attribute specified}}
    spv.EntryPoint "GLCompute" "do_nothing"
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // expected-error @+1 {{function 'do_something' not found in 'spv.module'}}
    spv.EntryPoint "GLCompute" @do_something
 }
 
 /// TODO(ravishankarm) : Add a test that verifies an error is thrown
 /// when interface entries of EntryPointOp are not
 /// spv.Variables. There is currently no other op that has a spv.ptr
 /// return type
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      // expected-error @+1 {{op must appear in a module-like op's block}}
      spv.EntryPoint "GLCompute" @do_something
    }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // expected-error @+1 {{duplicate of a previous EntryPointOp}}
    spv.EntryPoint "GLCompute" @do_nothing
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-   func @do_nothing() -> () {
+   spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // expected-error @+1 {{custom op 'spv.EntryPoint' invalid execution_model attribute specification: "ContractionOff"}}
    spv.EntryPoint "ContractionOff" @do_nothing
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
+// spv.func
+//===----------------------------------------------------------------------===//
+
+// CHECK: spv.func @foo() "None"
+spv.func @foo() "None"
+
+// CHECK: spv.func @bar(%{{.+}}: i32) -> i32 "Inline|Pure" {
+spv.func @bar(%arg: i32) -> (i32) "Inline|Pure" {
+  // CHECK-NEXT: spv.
+  spv.ReturnValue %arg: i32
+// CHECK-NEXT: }
+}
+
+// CHECK: spv.func @baz(%{{.+}}: i32) "DontInline" attributes {additional_stuff = 64 : i64}
+spv.func @baz(%arg: i32) "DontInline" attributes {
+  additional_stuff = 64
+} { spv.Return }
+
+// -----
+
+// expected-error @+1 {{expected function_control attribute specified as string}}
+spv.func @missing_function_control() { spv.Return }
+
+// -----
+
+// expected-error @+1 {{cannot have more than one result}}
+spv.func @cannot_have_more_than_one_result(%arg: i32) -> (i32, i32) "None"
+
+// -----
+
+// expected-error @+1 {{expected SSA identifier}}
+spv.func @cannot_have_variadic_arguments(%arg: i32, ...) "None"
+
+// -----
+
+// Nested function
+spv.module "Logical" "GLSL450" {
+  spv.func @outer_func() -> () "None" {
+    // expected-error @+1 {{must appear in a module-like op's block}}
+    spv.func @inner_func() -> () "None" {
+      spv.Return
+    }
+    spv.Return
+  }
+}
+
+// -----
+
+//===----------------------------------------------------------------------===//
 // spv.globalVariable
 //===----------------------------------------------------------------------===//
 
 spv.module "Logical" "GLSL450" {
   // CHECK: spv.globalVariable @var0 : !spv.ptr<f32, Input>
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
 }
 

 spv.module "Logical" "GLSL450" {
   // expected-error @+1 {{storage class cannot be 'Generic'}}
   spv.globalVariable @var0 : !spv.ptr<f32, Generic>
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @foo() {
+  spv.func @foo() "None" {
     // expected-error @+1 {{op must appear in a module-like op's block}}
     spv.globalVariable @var0 : !spv.ptr<f32, Input>
     spv.Return
   }
 }
 
 // -----
 

 // CHECK: spv.module
 spv.module "Logical" "GLSL450" {
   spv._module_end
 }
 
 // Module with function
 // CHECK: spv.module
 spv.module "Logical" "GLSL450" {
-  func @do_nothing() -> () {
+  spv.func @do_nothing() -> () "None" {
     spv.Return
   }
 }
 
 // -----
 
 // Missing addressing model
 // expected-error@+1 {{custom op 'spv.module' expected addressing_model attribute specified as string}}

 // expected-error@+2 {{expects regions to end with 'spv._module_end'}}
 // expected-note@+1 {{in custom textual format, the absence of terminator implies 'spv._module_end'}}
 "spv.module"() ({
   %0 = spv.constant true
 }) {addressing_model = 0 : i32, memory_model = 1 : i32} : () -> ()
 
 // -----
 
-// Use non SPIR-V op inside.module
+// Use non SPIR-V op inside module
 spv.module "Logical" "GLSL450" {
-  // expected-error @+1 {{'spv.module' can only contain func and spv.* ops}}
+  // expected-error @+1 {{'spv.module' can only contain spv.* ops}}
   "dialect.op"() : () -> ()
 }
 
 // -----
 
 // Use non SPIR-V op inside function
 spv.module "Logical" "GLSL450" {
-  func @do_nothing() -> () {
+  spv.func @do_nothing() -> () "None" {
     // expected-error @+1 {{functions in 'spv.module' can only contain spv.* ops}}
     "dialect.op"() : () -> ()
   }
 }
 
 // -----
 
 // Use external function
 spv.module "Logical" "GLSL450" {
   // expected-error @+1 {{'spv.module' cannot contain external functions}}
-  func @extern() -> ()
-}
-
-// -----
-
-// Module with nested function
-spv.module "Logical" "GLSL450" {
-  func @outer_func() -> () {
-    // expected-error @+1 {{'spv.module' cannot contain nested functions}}
-    func @inner_func() -> () {
-      spv.Return
-    }
-    spv.Return
-  }
+  spv.func @extern() -> () "None"
 }
 
 // -----
 
 // expected-error @+1 {{uses unknown capability: MyAwesomeCapability}}
 spv.module "Logical" "GLSL450" {
 } attributes {
   capabilities = ["MyAwesomeCapability"]

 //===----------------------------------------------------------------------===//
 
 spv.module "Logical" "GLSL450" {
   spv.specConstant @sc1 = false
   spv.specConstant @sc2 = 42 : i64
   spv.specConstant @sc3 = 1.5 : f32
 
   // CHECK-LABEL: @reference
-  func @reference() -> i1 {
+  spv.func @reference() -> i1 "None" {
     // CHECK: spv._reference_of @sc1 : i1
     %0 = spv._reference_of @sc1 : i1
     spv.ReturnValue %0 : i1
   }
 
   // CHECK-LABEL: @initialize
-  func @initialize() -> i64 {
+  spv.func @initialize() -> i64 "None" {
     // CHECK: spv._reference_of @sc2 : i64
     %0 = spv._reference_of @sc2 : i64
     %1 = spv.Variable init(%0) : !spv.ptr<i64, Function>
     %2 = spv.Load "Function" %1 : i64
     spv.ReturnValue %2 : i64
   }
 
   // CHECK-LABEL: @compute
-  func @compute() -> f32 {
+  spv.func @compute() -> f32 "None" {
     // CHECK: spv._reference_of @sc3 : f32
     %0 = spv._reference_of @sc3 : f32
     %1 = spv.constant 6.0 : f32
     %2 = spv.FAdd %0, %1 : f32
     spv.ReturnValue %2 : f32
   }
 }
 
 // -----
 
 // Allow taking reference of spec constant in other module-like ops
 spv.specConstant @sc = 5 : i32
 func @reference_of() {
   // CHECK: spv._reference_of @sc
   %0 = spv._reference_of @sc : i32
   return
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
     // expected-error @+1 {{expected spv.specConstant symbol}}
     %0 = spv._reference_of @sc : i32
     spv.Return
   }
 }
 
 // -----
 
 spv.module "Logical" "GLSL450" {
   spv.specConstant @sc = 42 : i32
-  func @foo() -> () {
+  spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}}
     %0 = spv._reference_of @sc : f32
     spv.Return
   }
 }
 
 // -----