[MLIR][SPIRV] Rename `spv.module` to `spv.mlir.module`.

[ergawy]

Summary

To unify the naming scheme across all ops in the SPIR-V dialect, we are
moving from spv.camelCase to spv.CamelCase everywhere. For ops that
don't have a SPIR-V spec counterpart, we use spv.mlir.snake_case.

[antiagainst]

Actually we can probably keep spv.module as-is. Similarly to spv.func, they maintain consistency with other module/func op names. What do you think?

[ergawy]

In D98017#2608850, @antiagainst wrote:

Actually we can probably keep spv.module as-is. Similarly to spv.func, they maintain consistency with other module/func op names. What do you think?

Yes, that makes sense. I will abandon then.

mlir/docs/Dialects/SPIR-V.md

*   All instructions in an extended instruction set are further qualified with
     GLSL extended instruction set have the prefix of `spv.GLSL.`.
 *   Ops that directly mirror instructions in the specification have `CamelCase`
     names that are the same as the instruction opnames (without the `Op`
     prefix). For example, `spv.FMul` is a direct mirror of `OpFMul` in the
     specification. Such an op will be serialized into and deserialized from one
     SPIR-V instruction.
 *   Ops with `snake_case` names are those that have different representation
     from corresponding instructions (or concepts) in the specification. These
-    ops are mostly for defining the SPIR-V structure. For example, `spv.module`
+    ops are mostly for defining the SPIR-V structure. For example, `spv.mlir.module`
     and `spv.Constant`. They may correspond to one or more instructions during
     (de)serialization.
 *   Ops with `mlir.snake_case` names are those that have no corresponding
     instructions (or concepts) in the binary format. They are introduced to
     satisfy MLIR structural requirements. For example, `spv.mlir.endmodule` and
     `spv.mlir.merge`. They map to no instructions during (de)serialization.
 
 (TODO: consider merging the last two cases and adopting `spv.mlir.` prefix for
 them.)
 
 ## Module
 
-A SPIR-V module is defined via the `spv.module` op, which has one region that
+A SPIR-V module is defined via the `spv.mlir.module` op, which has one region that
 contains one block. Model-level instructions, including function definitions,
 are all placed inside the block. Functions are defined using the builtin `func`
 op.
 
-We choose to model a SPIR-V module with a dedicated `spv.module` op based on the
+We choose to model a SPIR-V module with a dedicated `spv.mlir.module` op based on the
 following considerations:
 
 *   It maps cleanly to a SPIR-V module in the specification.
 *   We can enforce SPIR-V specific verification that is suitable to be performed
     at the module-level.
 *   We can attach additional model-level attributes.
 *   We can control custom assembly form.
 
-The `spv.module` op's region cannot capture SSA values from outside, neither
-implicitly nor explicitly. The `spv.module` op's region is closed as to what ops
+The `spv.mlir.module` op's region cannot capture SSA values from outside, neither
+implicitly nor explicitly. The `spv.mlir.module` op's region is closed as to what ops
 can appear inside: apart from the builtin `func` op, it can only contain ops
-from the SPIR-V dialect. The `spv.module` op's verifier enforces this rule. This
-meaningfully guarantees that a `spv.module` can be the entry point and boundary
+from the SPIR-V dialect. The `spv.mlir.module` op's verifier enforces this rule. This
+meaningfully guarantees that a `spv.mlir.module` can be the entry point and boundary
 for serialization.
 
 ### Module-level operations
 
 SPIR-V binary format defines the following [sections][SpirvLogicalLayout]:
 
 1.  Capabilities required by the module.
 1.  Extensions required by the module.

 resource variables to interact with the execution environment.
 
 Compared to the binary format, we adjust how these module-level SPIR-V
 instructions are represented in the SPIR-V dialect:
 
 #### Use MLIR attributes for metadata
 
 *   Requirements for capabilities, extensions, extended instruction sets,
-    addressing model, and memory model are conveyed using `spv.module`
+    addressing model, and memory model are conveyed using `spv.mlir.module`
     attributes. This is considered better because these information are for the
     execution environment. It's easier to probe them if on the module op itself.
 *   Annotations/decoration instructions are "folded" into the instructions they
     decorate and represented as attributes on those ops. This eliminates
     potential forward references of SSA values, improves IR readability, and
     makes querying the annotations more direct. More discussions can be found in
     the [`Decorations`](#decorations) section.
 
 #### Model types with MLIR custom types
 
 *   Types are represented using MLIR builtin types and SPIR-V dialect specific
     types. There are no type declaration ops in the SPIR-V dialect. More
     discussions can be found in the [Types](#types) section later.
 
 #### Unify and localize constants
 
 *   Various normal constant instructions are represented by the same
     `spv.Constant` op. Those instructions are just for constants of different
     types; using one op to represent them reduces IR verbosity and makes
     transformations less tedious.
-*   Normal constants are not placed in `spv.module`'s region; they are localized
+*   Normal constants are not placed in `spv.mlir.module`'s region; they are localized
     into functions. This is to make functions in the SPIR-V dialect to be
     isolated and explicit capturing. Constants are cheap to duplicate given
     attributes are made unique in `MLIRContext`.
 
 #### Adopt symbol-based global variables and specialization constant
 
 *   Global variables are defined with the `spv.GlobalVariable` op. They do not
     generate SSA values. Instead they have symbols and should be referenced via
*   In SPIR-V specification, functions support implicit capturing: they can
     defined with `func` op, which disallows implicit capturing. This is more
     friendly to compiler analyses and transformations. More discussions can be
     found in the [Function](#function) section later.
 
 #### Model entry points and execution models as normal ops
 
 *   A SPIR-V module can have multiple entry points. And these entry points refer
     to the function and interface variables. It’s not suitable to model them as
-    `spv.module` op attributes. We can model them as normal ops of using symbol
+    `spv.mlir.module` op attributes. We can model them as normal ops of using symbol
     references.
 *   Similarly for execution modes, which are coupled with entry points, we can
-    model them as normal ops in `spv.module`'s region.
+    model them as normal ops in `spv.mlir.module`'s region.
 
 ## Decorations
 
 Annotations/decorations provide additional information on result ids. In SPIR-V,
 all instructions can generate result ids, including value-computing and
 type-defining ones.
 
 For decorations on value result ids, we can just have a corresponding attribute

 
 ```mlir
 func @f(%arg: i32) -> i32 {
   "spv.ReturnValue"(%arg) : (i32) -> (i32)
 }
 ```
 
 A SPIR-V function can have at most one result. It cannot contain nested
-functions or non-SPIR-V operations. `spv.module` verifies these requirements.
+functions or non-SPIR-V operations. `spv.mlir.module` verifies these requirements.
 
 A major difference between the SPIR-V dialect and the SPIR-V specification for
 functions is that the former are isolated and require explicit capturing, while
 the latter allows implicit capturing. In SPIR-V specification, functions can
 refer to SSA values (generated by constants, global variables, etc.) defined in
 modules. The SPIR-V dialect adjusted how constants and global variables are
 modeled to enable isolated functions. Isolated functions are more friendly to
 compiler analyses and transformations. This also enables the SPIR-V dialect to

 binary module; so serialization is not designed to be a general tool for
 investigating the SPIR-V binary module and does not guarantee roundtrip
 equivalence (at least for now). For the latter, please use the
 assembler/disassembler in the [SPIRV-Tools][SpirvTools] project.
 
 A few transformations are performed in the process of serialization because of
 the representational differences between SPIR-V dialect and binary format:
 
-*   Attributes on `spv.module` are emitted as their corresponding SPIR-V
+*   Attributes on `spv.mlir.module` are emitted as their corresponding SPIR-V
     instructions.
 *   Types are serialized into `OpType*` instructions in the SPIR-V binary module
     section for types, constants, and global variables.
 *   `spv.Constant`s are unified and placed in the SPIR-V binary module section
     for types, constants, and global variables.
 *   Attributes on ops, if not part of the op's binary encoding, are emitted as
     `OpDecorate*` instructions in the SPIR-V binary module section for
     decorations.
 *   `spv.selection`s and `spv.loop`s are emitted as basic blocks with `Op*Merge`
     instructions in the header block as required by the binary format.
 *   Block arguments are materialized as `OpPhi` instructions at the beginning of
     the corresponding blocks.
 
 Similarly, a few transformations are performed during deserialization:
 
 *   Instructions for execution environment requirements (extensions,
     capabilities, extended instruction sets, etc.) will be placed as attributes
-    on `spv.module`.
+    on `spv.mlir.module`.
 *   `OpType*` instructions will be converted into proper `mlir::Type`s.
 *   `OpConstant*` instructions are materialized as `spv.Constant` at each use
     site.
 *   `OpVariable` instructions will be converted to `spv.GlobalVariable` ops if
     in module-level; otherwise they will be converted into `spv.Variable` ops.
 *   Every use of a module-level `OpVariable` instruction will materialize a
     `spv.mlir.addressof` op to turn the symbol of the corresponding
     `spv.GlobalVariable` into an SSA value.

 access to an instance of `mlir::SPIRVTypeLowering` class.
 
 ### Utility functions for lowering
 
 #### Setting shader interface
 
 The method `mlir::spirv::setABIAttrs` allows setting the [shader interface
 attributes](#shader-interface-abi) for a function that is to be an entry
-point function within the `spv.module` on lowering. A later pass
+point function within the `spv.mlir.module` on lowering. A later pass
 `mlir::spirv::LowerABIAttributesPass` uses this information to lower the entry
 point function and its ABI consistent with the Vulkan validation
 rules. Specifically,
 
 *   Creates `spv.GlobalVariable`s for the arguments, and replaces all uses of
     the argument with this variable. The SSA value used for replacement is
     obtained using the `spv.mlir.addressof` operation.
 *   Adds the `spv.EntryPoint` and `spv.ExecutionMode` operations into the
-    `spv.module` for the entry function.
+    `spv.mlir.module` for the entry function.
 
 #### Setting layout for shader interface variables
 
 SPIR-V validation rules for shaders require composite objects to be explicitly
 laid out. If a `spv.GlobalVariable` is not explicitly laid out, the utility
 method `mlir::spirv::decorateType` implements a layout consistent with
 the [Vulkan shader requirements][VulkanShaderInterface].
 
 #### Creating builtin variables
 
 In SPIR-V dialect, builtins are represented using `spv.GlobalVariable`s, with
 `spv.mlir.addressof` used to get a handle to the builtin as an SSA value.  The
 method `mlir::spirv::getBuiltinVariableValue` creates a `spv.GlobalVariable` for
-the builtin in the current `spv.module` if it does not exist already, and
+the builtin in the current `spv.mlir.module` if it does not exist already, and
 returns an SSA value generated from an `spv.mlir.addressof` operation.
 
 ### Current conversions to SPIR-V
 
 Using the above infrastructure, conversions are implemented from
 
 *   [Standard Dialect][MlirStandardDialect] : Only arithmetic and logical
     operations conversions are implemented.
-*   [GPU Dialect][MlirGpuDialect] : A gpu.module is converted to a `spv.module`.
+*   [GPU Dialect][MlirGpuDialect] : A gpu.module is converted to a `spv.mlir.module`.
     A gpu.function within this module is lowered as an entry function.
 
 ## Code organization
 
 We aim to provide multiple libraries with clear dependencies for SPIR-V related
 functionalities in MLIR so developers can just choose the needed components
 without pulling in the whole world.
 

 
 To lower an operation into SPIR-V dialect, implement a [conversion
 pattern][MlirDialectConversionRewritePattern]. If the conversion requires type
 conversion as well, the pattern must inherit from the
 `mlir::spirv::SPIRVOpLowering` class to get access to
 `mlir::spirv::SPIRVTypeConverter`.  If the operation has a region, [signature
 conversion][MlirDialectConversionSignatureConversion] might be needed as well.
 
-**Note**: The current validation rules of `spv.module` require that all
+**Note**: The current validation rules of `spv.mlir.module` require that all
 operations contained within its region are valid operations in the SPIR-V
 dialect.
 
 ## Operation definitions
 
 [include "Dialects/SPIRVOps.md"]
 
 [Spirv]: https://www.khronos.org/registry/spir-v/

mlir/docs/PassManagement.md

 `PassManager` acts as an `OpPassManager`. Nesting in this sense, corresponds to
 the [structural](Tutorials/UnderstandingTheIRStructure.md) nesting within
 [Regions](LangRef.md#regions) of the IR.
 
 For example, the following `.mlir`:
 
 ```
 module {
-  spv.module "Logical" "GLSL450" {
+  spv.mlir.module "Logical" "GLSL450" {
     func @foo() {
       ...
     }
   }
 }
 ```
 
 Has the nesting structure of:
 
 ```
 `module`
-  `spv.module`
+  `spv.mlir.module`
     `function`
 ```
 
 Below is an example of constructing a pipeline that operates on the above
 structure:
 
 ```c++
 // Create a top-level `PassManager` class. If an operation type is not

mlir/docs/SPIRVToLLVMDialectConversion.md

 
 In SPIR-V dialect, the global variable returns a pointer, whereas in LLVM
 dialect the global holds an actual value. This difference is handled by
 `spv.mlir.addressof` and `llvm.mlir.addressof` ops that both return a pointer and
 are used to reference the global.
 
 ```mlir
 // Original SPIR-V module
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @struct : !spv.ptr<!spv.struct<f32, !spv.array<10xf32>>, Private>
   spv.func @func() -> () "None" {
     %0 = spv.mlir.addressof @struct : !spv.ptr<!spv.struct<f32, !spv.array<10xf32>>, Private>
     spv.Return
   }
 }
 
 // Converted result

 ### `spv.Return` and `spv.ReturnValue`
 
 In LLVM IR, functions may return either 1 or 0 value. Hence, we map both ops to
 `llvm.return` with or without a return value.
 
 ## Module ops
 
 Module in SPIR-V has one region that contains one block. It is defined via
-`spv.module` op that also takes a range of attributes:
+`spv.mlir.module` op that also takes a range of attributes:
 
 *   Addressing model
 *   Memory model
 *   Version-Capability-Extension attribute
 
-`spv.module` is converted into `ModuleOp`. This plays a role of enclosing scope
+`spv.mlir.module` is converted into `ModuleOp`. This plays a role of enclosing scope
 to LLVM ops. At the moment, SPIR-V module attributes are ignored.
 
 `spv.mlir.endmodule` is mapped to an equivalent terminator `ModuleTerminatorOp`.
 
 ## `mlir-spirv-cpu-runner`
 
 `mlir-spirv-cpu-runner` allows to execute `gpu` dialect kernel on the CPU via
 SPIR-V to LLVM dialect conversion. Currently, only single-threaded kernel is
"gpu.launch_func"(/*grid dimensions*/, %buffer) {
     kernel = @foo::bar
   }
 }
 ```
 
 Lowering `gpu` dialect to SPIR-V dialect results in
 
 ```mlir
-spv.module @__spv__foo /*VCE triple and other metadata here*/ {
+spv.mlir.module @__spv__foo /*VCE triple and other metadata here*/ {
   spv.GlobalVariable @__spv__foo_arg bind(0,0) : ...
   spv.func @bar() {
     // Kernel code.
   }
   spv.EntryPoint @bar, ...
 }
 
 func @main() {
   // Fill the buffer with some data.
   %buffer = alloc : memref<8xi32>
   %data = ...
   call fillBuffer(%buffer, %data)
 
   "gpu.launch_func"(/*grid dimensions*/, %buffer) {
     kernel = @foo::bar
   }
 }
 ```
 
 Then, the lowering from standard dialect to LLVM dialect is applied to the host
 code.
 
 ```mlir
-spv.module @__spv__foo /*VCE triple and other metadata here*/ {
+spv.mlir.module @__spv__foo /*VCE triple and other metadata here*/ {
   spv.GlobalVariable @__spv__foo_arg bind(0,0) : ...
   spv.func @bar() {
     // Kernel code.
   }
   spv.EntryPoint @bar, ...
 }
 
 // Kernel function declaration.

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

let extraClassDeclaration = [{
     ::mlir::spirv::StorageClass storageClass() {
       return this->type().cast<::mlir::spirv::PointerType>().getStorageClass();
     }
   }];
 }
 
 // -----
 
-def SPV_ModuleOp : SPV_Op<"module",
+def SPV_ModuleOp : SPV_Op<"mlir.module",
                           [IsolatedFromAbove,
                            SingleBlockImplicitTerminator<"ModuleEndOp">,
                            SymbolTable, Symbol]> {
   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,
let description = [{
     This op has only one region, which only contains one block. The block
     must be terminated via the `spv.mlir.endmodule` op.
 
     <!-- End of AutoGen section -->
 
     ```
     addressing-model ::= `Logical` | `Physical32` | `Physical64` | ...
     memory-model ::= `Simple` | `GLSL450` | `OpenCL` | `Vulkan` | ...
-    spv-module-op ::= `spv.module` addressing-model memory-model
+    spv-module-op ::= `spv.mlir.module` addressing-model memory-model
                       (requires  spirv-vce-attribute)?
                       (`attributes` attribute-dict)?
                       region
     ```
 
     #### Example:
 
     ```mlir
-    spv.module Logical GLSL450  {}
+    spv.mlir.module Logical GLSL450  {}
 
-    spv.module Logical Vulkan
+    spv.mlir.module Logical Vulkan
         requires #spv.vce<v1.0, [Shader], [SPV_KHR_vulkan_memory_model]>
         attributes { some_additional_attr = ... } {
       spv.func @do_nothing() -> () {
         spv.Return
       }
     }
     ```
   }];

 }
 
 // -----
 
 def SPV_ModuleEndOp : SPV_Op<"mlir.endmodule", [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 terminates the only block inside a `spv.mlir.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);

mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h

 /// Creates a module pass that converts composite types used by objects in the
 /// StorageBuffer, PhysicalStorageBuffer, Uniform, and PushConstant storage
 /// classes with layout information.
 /// Right now this pass only supports Vulkan layout rules.
 std::unique_ptr<OperationPass<mlir::ModuleOp>>
 createDecorateSPIRVCompositeTypeLayoutPass();
 
 /// Creates an operation pass that deduces and attaches the minimal version/
-/// capabilities/extensions requirements for spv.module ops.
-/// For each spv.module op, this pass requires a `spv.target_env` attribute on
-/// it or an enclosing module-like op to drive the deduction. The reason is
+/// capabilities/extensions requirements for spv.mlir.module ops.
+/// For each spv.mlir.module op, this pass requires a `spv.target_env` attribute
+/// on it or an enclosing module-like op to drive the deduction. The reason is
 /// that an op can be enabled by multiple extensions/capabilities. So we need
 /// to know which one to pick. `spv.target_env` gives the hard limit as for
 /// what the target environment can support; this pass deduces what are
-/// actually needed for a specific spv.module op.
+/// actually needed for a specific spv.mlir.module op.
 std::unique_ptr<OperationPass<spirv::ModuleOp>>
 createUpdateVersionCapabilityExtensionPass();
 
 /// Creates an operation pass that lowers the ABI attributes specified during
 /// SPIR-V Lowering. Specifically,
 /// 1. Creates the global variables for arguments of entry point function using
 ///    the specification in the `spv.interface_var_abi` attribute for each
 ///    argument.

mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td

 def SPIRVRewriteInsertsPass : Pass<"spirv-rewrite-inserts", "spirv::ModuleOp"> {
   let summary = "Rewrite sequential chains of spv.CompositeInsert operations into "
                 "spv.CompositeConstruct operations";
   let constructor = "mlir::spirv::createRewriteInsertsPass()";
 }
 
 def SPIRVUpdateVCE : Pass<"spirv-update-vce", "spirv::ModuleOp"> {
   let summary = "Deduce and attach minimal (version, capabilities, extensions) "
-                "requirements to spv.module ops";
+                "requirements to spv.mlir.module ops";
   let constructor = "mlir::spirv::createUpdateVersionCapabilityExtensionPass()";
 }
 
 #endif // MLIR_DIALECT_SPIRV_TRANSFORMS_PASSES

mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp

 public:
   using OpConversionPattern<gpu::BlockDimOp>::OpConversionPattern;
 
   LogicalResult
   matchAndRewrite(gpu::BlockDimOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
-/// Pattern to convert a kernel function in GPU dialect within a spv.module.
+/// Pattern to convert a kernel function in GPU dialect within a
+/// spv.mlir.module.
 class GPUFuncOpConversion final : public OpConversionPattern<gpu::GPUFuncOp> {
 public:
   using OpConversionPattern<gpu::GPUFuncOp>::OpConversionPattern;
 
   LogicalResult
   matchAndRewrite(gpu::GPUFuncOp funcOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 
 private:
   SmallVector<int32_t, 3> workGroupSizeAsInt32;
 };
 
-/// Pattern to convert a gpu.module to a spv.module.
+/// Pattern to convert a gpu.module to a spv.mlir.module.
 class GPUModuleConversion final : public OpConversionPattern<gpu::GPUModuleOp> {
 public:
   using OpConversionPattern<gpu::GPUModuleOp>::OpConversionPattern;
 
   LogicalResult
   matchAndRewrite(gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
LogicalResult GPUModuleConversion::matchAndRewrite(
   auto spvModule = rewriter.create<spirv::ModuleOp>(
       moduleOp.getLoc(), addressingModel, memoryModel.getValue(),
       StringRef(spvModuleName));
 
   // Move the region from the module op into the SPIR-V module.
   Region &spvModuleRegion = spvModule.body();
   rewriter.inlineRegionBefore(moduleOp.body(), spvModuleRegion,
                               spvModuleRegion.begin());
-  // The spv.module build method adds a block with a terminator. Remove that
-  // block. The terminator of the module op in the remaining block will be
+  // The spv.mlir.module build method adds a block with a terminator. Remove
+  // that block. The terminator of the module op in the remaining block will be
   // legalized later.
   rewriter.eraseBlock(&spvModuleRegion.back());
   rewriter.eraseOp(moduleOp);
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // GPU return inside kernel functions to SPIR-V return.

mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp

 //===- GPUToSPIRVPass.cpp - GPU to SPIR-V Passes --------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a pass to convert a kernel function in the GPU Dialect
-// into a spv.module operation.
+// into a spv.mlir.module operation.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h"
 
 #include "../PassDetail.h"
 #include "mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h"
 #include "mlir/Conversion/StandardToSPIRV/StandardToSPIRV.h"

mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp

 }
 
 LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::createBinaryShader(
     ModuleOp module, std::vector<char> &binaryShader) {
   bool done = false;
   SmallVector<uint32_t, 0> binary;
   for (auto spirvModule : module.getOps<spirv::ModuleOp>()) {
     if (done)
-      return spirvModule.emitError("should only contain one 'spv.module' op");
+      return spirvModule.emitError(
+          "should only contain one 'spv.mlir.module' op");
     done = true;
 
     if (failed(spirv::serialize(spirvModule, binary)))
       return failure();
   }
   binaryShader.resize(binary.size() * sizeof(uint32_t));
   std::memcpy(binaryShader.data(), reinterpret_cast<char *>(binary.data()),
               binaryShader.size());

mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp

void ConvertSPIRVToLLVMPass::runOnOperation() {
   populateSPIRVToLLVMModuleConversionPatterns(context, converter, patterns);
   populateSPIRVToLLVMConversionPatterns(context, converter, patterns);
   populateSPIRVToLLVMFunctionConversionPatterns(context, converter, patterns);
 
   ConversionTarget target(getContext());
   target.addIllegalDialect<spirv::SPIRVDialect>();
   target.addLegalDialect<LLVM::LLVMDialect>();
 
-  // Set `ModuleOp` and `ModuleTerminatorOp` as legal for `spv.module`
+  // Set `ModuleOp` and `ModuleTerminatorOp` as legal for `spv.mlir.module`
   // conversion.
   target.addLegalOp<ModuleOp>();
   target.addLegalOp<ModuleTerminatorOp>();
   if (failed(applyPartialConversion(module, target, std::move(patterns))))
     signalPassFailure();
 }
 
 std::unique_ptr<OperationPass<ModuleOp>> mlir::createConvertSPIRVToLLVMPass() {
   return std::make_unique<ConvertSPIRVToLLVMPass>();
 }

mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp

 // Note that DRR cannot be used for the patterns in this file: we may need to
 // convert type along the way, which requires ConversionPattern. DRR generates
 // normal RewritePattern.
 
 namespace {
 
 /// Converts an allocation operation to SPIR-V. Currently only supports lowering
 /// to Workgroup memory when the size is constant.  Note that this pattern needs
-/// to be applied in a pass that runs at least at spv.module scope since it wil
-/// ladd global variables into the spv.module.
+/// to be applied in a pass that runs at least at spv.mlir.module scope since it
+/// wil ladd global variables into the spv.mlir.module.
 class AllocOpPattern final : public OpConversionPattern<AllocOp> {
 public:
   using OpConversionPattern<AllocOp>::OpConversionPattern;
 
   LogicalResult
   matchAndRewrite(AllocOp operation, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     MemRefType allocType = operation.getType();

mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp

static LogicalResult verify(spirv::MergeOp mergeOp) {
   Block &parentLastBlock = mergeOp->getParentRegion()->back();
   if (mergeOp.getOperation() != parentLastBlock.getTerminator())
     return mergeOp.emitOpError(
         "can only be used in the last block of 'spv.selection' or 'spv.loop'");
   return success();
 }
 
 //===----------------------------------------------------------------------===//
-// spv.module
+// spv.mlir.module
 //===----------------------------------------------------------------------===//
 
 void spirv::ModuleOp::build(OpBuilder &builder, OperationState &state,
                             Optional<StringRef> name) {
   ensureTerminator(*state.addRegion(), builder, state.location);
   if (name) {
     state.attributes.append(mlir::SymbolTable::getSymbolAttrName(),
                             builder.getStringAttr(*name));
static LogicalResult verify(spirv::ModuleOp moduleOp) {
   auto &op = *moduleOp.getOperation();
   auto *dialect = op.getDialect();
   DenseMap<std::pair<spirv::FuncOp, spirv::ExecutionModel>, spirv::EntryPointOp>
       entryPoints;
   SymbolTable table(moduleOp);
 
   for (auto &op : moduleOp.getBlock()) {
     if (op.getDialect() != dialect)
-      return op.emitError("'spv.module' can only contain spv.* ops");
+      return op.emitError("'spv.mlir.module' can only contain spv.* ops");
 
     // 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<spirv::FuncOp>(entryPointOp.fn());
       if (!funcOp) {
         return entryPointOp.emitError("function '")
-               << entryPointOp.fn() << "' not found in 'spv.module'";
+               << entryPointOp.fn() << "' not found in 'spv.mlir.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 '")
if (auto entryPointOp = dyn_cast<spirv::EntryPointOp>(op)) {
           funcOp, entryPointOp.execution_model());
       auto entryPtIt = entryPoints.find(key);
       if (entryPtIt != entryPoints.end()) {
         return entryPointOp.emitError("duplicate of a previous EntryPointOp");
       }
       entryPoints[key] = entryPointOp;
     } else if (auto funcOp = dyn_cast<spirv::FuncOp>(op)) {
       if (funcOp.isExternal())
-        return op.emitError("'spv.module' cannot contain external functions");
+        return op.emitError(
+            "'spv.mlir.module' cannot contain external functions");
 
       // TODO: move this check to spv.func.
       for (auto &block : funcOp)
         for (auto &op : block) {
           if (op.getDialect() != dialect)
             return op.emitError(
-                "functions in 'spv.module' can only contain spv.* ops");
+                "functions in 'spv.mlir.module' can only contain spv.* ops");
         }
     }
   }
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//

mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp

case spirv::Opcode::OpName:
     return processName(operands);
   case spirv::Opcode::OpString:
     return processDebugString(operands);
   case spirv::Opcode::OpModuleProcessed:
   case spirv::Opcode::OpSource:
   case spirv::Opcode::OpSourceContinued:
   case spirv::Opcode::OpSourceExtension:
     // TODO: This is debug information embedded in the binary which should be
-    // translated into the spv.module.
+    // translated into the spv.mlir.module.
     return success();
   case spirv::Opcode::OpTypeVoid:
   case spirv::Opcode::OpTypeBool:
   case spirv::Opcode::OpTypeInt:
   case spirv::Opcode::OpTypeFloat:
   case spirv::Opcode::OpTypeVector:
   case spirv::Opcode::OpTypeMatrix:
   case spirv::Opcode::OpTypeArray:

mlir/lib/Target/SPIRV/Serialization/SerializeOps.cpp

Serializer::processOp<spirv::EntryPointOp>(spirv::EntryPointOp op) {
   // Add the name of the function.
   (void)spirv::encodeStringLiteralInto(operands, op.fn());
 
   // Add the interface values.
   if (auto interface = op.interface()) {
     for (auto var : interface.getValue()) {
       auto id = getVariableID(var.cast<FlatSymbolRefAttr>().getValue());
       if (!id) {
-        return op.emitError("referencing undefined global variable."
-                            "spv.EntryPoint is at the end of spv.module. All "
-                            "referenced variables should already be defined");
+        return op.emitError(
+            "referencing undefined global variable."
+            "spv.EntryPoint is at the end of spv.mlir.module. All "
       }
       operands.push_back(id);
     }
   }
   return encodeInstructionInto(entryPoints, spirv::Opcode::OpEntryPoint,
                                operands);
 }
 

mlir/lib/Target/SPIRV/TranslateRegistration.cpp

if (!module)
     return failure();
 
   SmallVector<uint32_t, 0> binary;
 
   SmallVector<spirv::ModuleOp, 1> spirvModules;
   module.walk([&](spirv::ModuleOp op) { spirvModules.push_back(op); });
 
   if (spirvModules.empty())
-    return module.emitError("found no 'spv.module' op");
+    return module.emitError("found no 'spv.mlir.module' op");
 
   if (spirvModules.size() != 1)
-    return module.emitError("found more than one 'spv.module' op");
+    return module.emitError("found more than one 'spv.mlir.module' op");
 
   if (failed(
           spirv::serialize(spirvModules[0], binary, /*emitDebuginfo=*/false)))
     return failure();
 
   output.write(reinterpret_cast<char *>(binary.data()),
                binary.size() * sizeof(uint32_t));
 

 
 static LogicalResult roundTripModule(ModuleOp srcModule, bool emitDebugInfo,
                                      raw_ostream &output) {
   SmallVector<uint32_t, 0> binary;
   MLIRContext *context = srcModule.getContext();
   auto spirvModules = srcModule.getOps<spirv::ModuleOp>();
 
   if (spirvModules.begin() == spirvModules.end())
-    return srcModule.emitError("found no 'spv.module' op");
+    return srcModule.emitError("found no 'spv.mlir.module' op");
 
   if (std::next(spirvModules.begin()) != spirvModules.end())
-    return srcModule.emitError("found more than one 'spv.module' op");
+    return srcModule.emitError("found more than one 'spv.mlir.module' op");
 
   if (failed(spirv::serialize(*spirvModules.begin(), binary, emitDebugInfo)))
     return failure();
 
   MLIRContext deserializationContext(context->getDialectRegistry());
   // TODO: we should only load the required dialects instead of all dialects.
   deserializationContext.loadAllAvailableDialects();
   // Then deserialize to get back a SPIR-V module.

mlir/test/Conversion/GPUToSPIRV/builtins.mlir

 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_y
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_y() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "y"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_z() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "z"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}} {
       // The constant value is obtained from the spv.entry_point_abi.
       // Note that this ignores the workgroup size specification in gpu.launch.
       // We may want to define gpu.workgroup_size and convert it to the entry
       // point ABI we want here.
       // CHECK: spv.Constant 32 : i32
       %0 = "gpu.block_dim"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_y
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_y() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // The constant value is obtained from the spv.entry_point_abi.
       // CHECK: spv.Constant 4 : i32
       %0 = "gpu.block_dim"() {dimension = "y"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_z() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // The constant value is obtained from the spv.entry_point_abi.
       // CHECK: spv.Constant 1 : i32
       %0 = "gpu.block_dim"() {dimension = "z"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_local_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   // CHECK: spv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
   gpu.module @kernels {
     gpu.func @builtin_local_id_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[LOCALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.thread_id"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_num_workgroups_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   // CHECK: spv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
   gpu.module @kernels {
     gpu.func @builtin_num_workgroups_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMWORKGROUPS]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.grid_dim"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   // CHECK: spv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId")
   gpu.module @kernels {
     gpu.func @builtin_subgroup_id() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPID]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_id : index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   // CHECK: spv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups")
   gpu.module @kernels {
     gpu.func @builtin_num_subgroups() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMSUBGROUPS]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.num_subgroups : index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
-  // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL:  spv.mlir.module @{{.*}} Logical GLSL450
   // CHECK: spv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize")
   gpu.module @kernels {
     gpu.func @builtin_subgroup_size() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPSIZE]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_size : index
       gpu.return
     }
   }
 }

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

func @load_store(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>) {
     %c1_2 = constant 1 : index
     gpu.launch_func @kernels::@load_store_kernel
         blocks in (%0, %c1_2, %c1_2) threads in (%1, %c1_2, %c1_2)
         args(%arg0 : memref<12x4xf32>, %arg1 : memref<12x4xf32>, %arg2 : memref<12x4xf32>, 
              %c0 : index, %c0_0 : index, %c1 : index, %c1_1 : index)
     return
   }
 
-  // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
+  // CHECK-LABEL: spv.mlir.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:    spv.func @load_store_kernel
     // CHECK-SAME: %[[ARG0:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 0)>}
     // CHECK-SAME: %[[ARG1:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
     // CHECK-SAME: %[[ARG2:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 2)>}

mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, {}>
 } {
   gpu.module @kernels {
-    // CHECK-LABEL: spv.module @{{.*}} Physical64 OpenCL
+    // CHECK-LABEL: spv.mlir.module @{{.*}} Physical64 OpenCL
     //       CHECK:   spv.func
     //  CHECK-SAME:     {{%.*}}: f32
     //   CHECK-NOT:     spv.interface_var_abi
     //  CHECK-SAME:     {{%.*}}: !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, CrossWorkgroup>
     //   CHECK-NOT:     spv.interface_var_abi
     //  CHECK-SAME:     spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
     gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, 11>) kernel
         attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {

mlir/test/Conversion/GPUToSPIRV/simple.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -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:       spv.mlir.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spv.func @basic_module_structure
     // CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>}
     // CHECK-SAME: {{%.*}}: !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
     // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
     gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32>) kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // CHECK: spv.Return
       gpu.return
func @main() {
     return
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
-    // CHECK:       spv.module @{{.*}} Logical GLSL450 {
+    // CHECK:       spv.mlir.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spv.func @basic_module_structure_preset_ABI
     // CHECK-SAME: {{%[a-zA-Z0-9_]*}}: f32
     // CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>
     // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
     // CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(3, 0)>
     // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
     gpu.func @basic_module_structure_preset_ABI(
       %arg0 : f32

mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir

 // RUN: mlir-opt %s -convert-gpu-launch-to-vulkan-launch | FileCheck %s
 
 // CHECK: %[[resource:.*]] = alloc() : memref<12xf32>
 // CHECK: %[[index:.*]] = constant 1 : index
 // CHECK: call @vulkanLaunch(%[[index]], %[[index]], %[[index]], %[[resource]]) {spirv_blob = "{{.*}}", spirv_entry_point = "kernel"}
 
 module attributes {gpu.container_module} {
-  spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
+  spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
     spv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
     spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
       %2 = spv.Constant 0 : i32
       %3 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
       %4 = spv.AccessChain %0[%2, %2] : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
       %5 = spv.Load "StorageBuffer" %4 : f32
       spv.Return

mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir

 // RUN: mlir-opt -convert-spirv-to-llvm -split-input-file -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spv.Branch
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @branch_without_arguments() -> () "None" {
     // CHECK: llvm.br ^bb1
     spv.Branch ^label
   // CHECK: ^bb1
   ^label:
     spv.Return
   }
 

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.BranchConditional
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @cond_branch_without_arguments() -> () "None" {
     // CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1
     %cond = spv.Constant true
     // CHECK: lvm.cond_br %[[COND]], ^bb1, ^bb2
     spv.BranchConditional %cond, ^true, ^false
     // CHECK: ^bb1:
   ^true:
     spv.Return

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.loop
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK-LABEL: @infinite_loop
   spv.func @infinite_loop(%count : i32) -> () "None" {
     // CHECK:   llvm.br ^[[BB1:.*]]
     // CHECK: ^[[BB1]]:
     // CHECK:   %[[COND:.*]] = llvm.mlir.constant(true) : i1
     // CHECK:   llvm.cond_br %[[COND]], ^[[BB2:.*]], ^[[BB4:.*]]
     // CHECK: ^[[BB2]]:
     // CHECK:   llvm.br ^[[BB3:.*]]

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.selection
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @selection_empty() -> () "None" {
     // CHECK: llvm.return
     spv.selection {
     }
     spv.Return
   }
 
   spv.func @selection_with_merge_block_only() -> () "None" {

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

 // RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s
   
 module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
 
   //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() : !llvm.struct<(array<6 x i32>)>
   //       CHECK: llvm.func @__spv__foo_bar()
 
-  //       CHECK: spv.module @__spv__foo
+  //       CHECK: spv.mlir.module @__spv__foo
   //       CHECK:   spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
   //       CHECK:   spv.func @__spv__foo_bar
   
   //       CHECK:   spv.EntryPoint "GLCompute" @__spv__foo_bar
   //       CHECK:   spv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1
 
   // CHECK-LABEL: @main
   //       CHECK:   %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr<i32>, ptr<i32>, i64, array<1 x i64>, array<1 x i64>)>
   //  CHECK-NEXT:   %[[DEST:.*]] = llvm.mlir.addressof @__spv__foo_bar_arg_0_descriptor_set0_binding0 : !llvm.ptr<struct<(array<6 x i32>)>>
   //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]], %{{.*}}) : (!llvm.ptr<struct<(array<6 x i32>)>>, !llvm.ptr<i32>, i64, i1) -> ()
   //  CHECK-NEXT:   llvm.call @__spv__foo_bar() : () -> ()
   //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr<i32>, !llvm.ptr<struct<(array<6 x i32>)>>, i64, i1) -> ()
 
-  spv.module @__spv__foo Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> {
+  spv.mlir.module @__spv__foo Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> {
     spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
     spv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spv.mlir.addressof @bar_arg_0 : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
       spv.Return
     }
     spv.EntryPoint "GLCompute" @bar
     spv.ExecutionMode @bar "LocalSize", 1, 1, 1
   }

mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir

spv.func @access_chain_array(%arg0 : i32) "None" {
   %2 = spv.Load "Function" %1 ["Volatile"] : !spv.array<4xf32>
   spv.Return
 }
 
 //===----------------------------------------------------------------------===//
 // spv.GlobalVariable and spv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: llvm.mlir.global external constant @var() : f32
   spv.GlobalVariable @var : !spv.ptr<f32, Input>
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   //       CHECK: llvm.mlir.global private @struct() : !llvm.struct<packed (f32, array<10 x f32>)>
   // CHECK-LABEL: @func
   //       CHECK:   llvm.mlir.addressof @struct : !llvm.ptr<struct<packed (f32, array<10 x f32>)>>
   spv.GlobalVariable @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
   spv.func @func() "None" {
     %0 = spv.mlir.addressof @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
     spv.Return
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   //       CHECK: llvm.mlir.global external @bar_descriptor_set0_binding0() : i32
   // CHECK-LABEL: @foo
   //       CHECK:   llvm.mlir.addressof @bar_descriptor_set0_binding0 : !llvm.ptr<i32>
   spv.GlobalVariable @bar bind(0, 0) : !spv.ptr<i32, StorageBuffer>
   spv.func @foo() "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }
 }
 
-spv.module @name Logical GLSL450 {
+spv.mlir.module @name Logical GLSL450 {
   //       CHECK: llvm.mlir.global external @name_bar_descriptor_set0_binding0() : i32
   // CHECK-LABEL: @foo
   //       CHECK:   llvm.mlir.addressof @name_bar_descriptor_set0_binding0 : !llvm.ptr<i32>
   spv.GlobalVariable @bar bind(0, 0) : !spv.ptr<i32, StorageBuffer>
   spv.func @foo() "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }

mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir

 // CHECK-NEXT:     %[[VAL:.*]] = llvm.mlir.constant(31 : i32) : i32
 // CHECK-NEXT:     %[[RET:.*]] = llvm.insertvalue %[[VAL]], %[[UNDEF]][0 : i32] : !llvm.struct<(i32)>
 // CHECK-NEXT:     llvm.return %[[RET]] : !llvm.struct<(i32)>
 // CHECK-NEXT:   }
 // CHECK-NEXT:   llvm.func @empty
 // CHECK-NEXT:     llvm.return
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
-spv.module Logical OpenCL {
+spv.mlir.module Logical OpenCL {
   spv.func @empty() "None" {
     spv.Return
   }
   spv.EntryPoint "Kernel" @empty
   spv.ExecutionMode @empty "ContractionOff"
 }
 
 //      CHECK: module {
 // CHECK-NEXT:   llvm.mlir.global external constant @{{.*}}() : !llvm.struct<(i32, array<3 x i32>)> {
 // CHECK-NEXT:     %[[UNDEF:.*]] = llvm.mlir.undef : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     %[[EM:.*]] = llvm.mlir.constant(18 : i32) : i32
 // CHECK-NEXT:     %[[T0:.*]] = llvm.insertvalue %[[EM]], %[[UNDEF]][0 : i32] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     %[[C0:.*]] = llvm.mlir.constant(32 : i32) : i32
 // CHECK-NEXT:     %[[T1:.*]] = llvm.insertvalue %[[C0]], %[[T0]][1 : i32, 0 : i32] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     %[[C1:.*]] = llvm.mlir.constant(1 : i32) : i32
 // CHECK-NEXT:     %[[T2:.*]] = llvm.insertvalue %[[C1]], %[[T1]][1 : i32, 1 : i32] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     %[[C2:.*]] = llvm.mlir.constant(1 : i32) : i32
 // CHECK-NEXT:     %[[RET:.*]] = llvm.insertvalue %[[C2]], %[[T2]][1 : i32, 2 : i32] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     llvm.return %[[RET]] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:   }
 // CHECK-NEXT:   llvm.func @bar
 // CHECK-NEXT:     llvm.return
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
-spv.module Logical OpenCL {
+spv.mlir.module Logical OpenCL {
   spv.func @bar() "None" {
     spv.Return
   }
   spv.EntryPoint "Kernel" @bar
   spv.ExecutionMode @bar "LocalSizeHint", 32, 1, 1
 }
 
 //===----------------------------------------------------------------------===//

mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir

 // RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
-// spv.module
+// spv.mlir.module
 //===----------------------------------------------------------------------===//
 
 // CHECK: module
-spv.module Logical GLSL450 {}
+spv.mlir.module Logical GLSL450 {}
 
 // CHECK: module @foo
-spv.module @foo Logical GLSL450 {}
+spv.mlir.module @foo Logical GLSL450 {}
 
 // CHECK: module
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> {}
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> {}
 
 // CHECK: module
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: }
   spv.mlir.endmodule
 }
 
 // CHECK: module
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK-LABEL: llvm.func @empty()
   spv.func @empty() -> () "None" {
     // CHECK: llvm.return
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/IR/availability.mlir

func @subgroup_ballot(%predicate: i1) -> vector<4xi32> {
   // CHECK: extensions: [ ]
   // CHECK: capabilities: [ [GroupNonUniformBallot] ]
   %0 = spv.GroupNonUniformBallot Workgroup %predicate : vector<4xi32>
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: module_logical_glsl450
 func @module_logical_glsl450() {
-  // CHECK: spv.module min version: v1.0
-  // CHECK: spv.module max version: v1.5
-  // CHECK: spv.module extensions: [ ]
-  // CHECK: spv.module capabilities: [ [Shader] ]
-  spv.module Logical GLSL450 { }
+  // CHECK: spv.mlir.module min version: v1.0
+  // CHECK: spv.mlir.module max version: v1.5
+  // CHECK: spv.mlir.module extensions: [ ]
+  // CHECK: spv.mlir.module capabilities: [ [Shader] ]
+  spv.mlir.module Logical GLSL450 { }
   return
 }
 
 // CHECK-LABEL: module_physical_storage_buffer64_vulkan
 func @module_physical_storage_buffer64_vulkan() {
-  // CHECK: spv.module min version: v1.0
-  // CHECK: spv.module max version: v1.5
-  // CHECK: spv.module extensions: [ [SPV_EXT_physical_storage_buffer, SPV_KHR_physical_storage_buffer] [SPV_KHR_vulkan_memory_model] ]
-  // CHECK: spv.module capabilities: [ [PhysicalStorageBufferAddresses] [VulkanMemoryModel] ]
-  spv.module PhysicalStorageBuffer64 Vulkan { }
+  // CHECK: spv.mlir.module min version: v1.0
+  // CHECK: spv.mlir.module max version: v1.5
+  // CHECK: spv.mlir.module extensions: [ [SPV_EXT_physical_storage_buffer, SPV_KHR_physical_storage_buffer] [SPV_KHR_vulkan_memory_model] ]
+  // CHECK: spv.mlir.module capabilities: [ [PhysicalStorageBufferAddresses] [VulkanMemoryModel] ]
+  spv.mlir.module PhysicalStorageBuffer64 Vulkan { }
   return
 }

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

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.FunctionCall
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   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

 func @caller() {
   // CHECK: spv.FunctionCall
   spv.FunctionCall @callee() : () -> ()
   spv.Return
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @f_invalid_result_type(%arg0 : i32, %arg1 : i32) -> () "None" {
     // expected-error @+1 {{result group starting at #0 requires 0 or 1 element, but found 2}}
     %0:2 = spv.FunctionCall @f_invalid_result_type(%arg0, %arg1) : (i32, i32) -> (i32, i32)
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   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 {
+spv.mlir.module Logical GLSL450 {
   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 {
+spv.mlir.module Logical GLSL450 {
   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 {
+spv.mlir.module Logical GLSL450 {
   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.ReturnValue %cst: i32
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   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.ReturnValue %0: i32
   }
 }
 
 // -----

 "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 {
+spv.mlir.module Logical GLSL450 {
   spv.func @work() -> (i32) "None" {
     // expected-error @+1 {{cannot be used in functions returning value}}
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.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.mlir.merge

 "foo.function"() ({
   %0 = spv.Constant true
   // expected-error @+1 {{op must appear in a function-like op's block}}
   spv.ReturnValue %0 : i1
 })  : () -> ()
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.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
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.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.mlir.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:

mlir/test/Dialect/SPIRV/IR/matrix-ops.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -verify-diagnostics %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @matrix_times_scalar
   spv.func @matrix_times_scalar(%arg0 : !spv.matrix<3 x vector<3xf32>>, %arg1 : f32) -> !spv.matrix<3 x vector<3xf32>> "None" {
     // CHECK: {{%.*}} = spv.MatrixTimesScalar {{%.*}}, {{%.*}} : !spv.matrix<3 x vector<3xf32>>, f32 -> !spv.matrix<3 x vector<3xf32>>
     %result = spv.MatrixTimesScalar %arg0, %arg1 : !spv.matrix<3 x vector<3xf32>>, f32 -> !spv.matrix<3 x vector<3xf32>>
     spv.ReturnValue %result : !spv.matrix<3 x vector<3xf32>>
   }
 
   // CHECK-LABEL: @matrix_transpose_1

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

func @aligned_load_incorrect_attributes() -> () {
   %0 = spv.Variable : !spv.ptr<f32, Function>
   // expected-error @+1 {{expected ']'}}
   %1 = spv.Load "Function" %0 ["Aligned", 4, 23] : f32
   return
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
   // CHECK_LABEL: @simple_load
   spv.func @simple_load() -> () "None" {
     // CHECK: spv.Load "Input" {{%.*}} : f32
     %0 = spv.mlir.addressof @var0 : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.Return
   }
func @aligned_store_incorrect_attributes(%arg0 : f32) -> () {
   %0 = spv.Variable : !spv.ptr<f32, Function>
   // expected-error @+1 {{expected ']'}}
   spv.Store  "Function" %0, %arg0 ["Aligned", 4, 23] : f32
   return
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
   spv.func @simple_store(%arg0 : f32) -> () "None" {
     %0 = spv.mlir.addressof @var0 : !spv.ptr<f32, Input>
     // CHECK: spv.Store  "Input" {{%.*}}, {{%.*}} : f32
     spv.Store  "Input" %0, %arg0 : f32
     spv.Return
   }
 }
func @variable_init_normal_constant() -> () {
   %0 = spv.Constant 4.0 : f32
   // CHECK: spv.Variable init(%0) : !spv.ptr<f32, Function>
   %1 = spv.Variable init(%0) : !spv.ptr<f32, Function>
   return
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @global : !spv.ptr<f32, Workgroup>
   spv.func @variable_init_global_variable() -> () "None" {
     %0 = spv.mlir.addressof @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
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc = 42 : i32
   // CHECK-LABEL: @variable_init_spec_constant
   spv.func @variable_init_spec_constant() -> () "None" {
     %0 = spv.mlir.referenceof @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/IR/structure-ops.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   spv.func @access_chain() -> () "None" {
     %0 = spv.Constant 1: i32
     // CHECK: [[VAR1:%.*]] = spv.mlir.addressof @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.mlir.addressof @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>, i32, i32
     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 @addressof() -> () {
   // CHECK: spv.mlir.addressof @var
   %1 = spv.mlir.addressof @var : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   return
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{expected spv.GlobalVariable symbol}}
     %0 = spv.mlir.addressof @var2 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced global variable's type}}
     %0 = spv.mlir.addressof @var1 : !spv.ptr<f32, Input>
   }
 }
 
 // -----

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.EntryPoint
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // CHECK: spv.EntryPoint "GLCompute" @do_nothing
    spv.EntryPoint "GLCompute" @do_nothing
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
    spv.GlobalVariable @var2 : !spv.ptr<f32, Input>
    spv.GlobalVariable @var3 : !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 {
+spv.mlir.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // expected-error @+1 {{invalid kind of attribute specified}}
    spv.EntryPoint "GLCompute" "do_nothing"
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
-   // expected-error @+1 {{function 'do_something' not found in 'spv.module'}}
+   // expected-error @+1 {{function 'do_something' not found in 'spv.mlir.module'}}
    spv.EntryPoint "GLCompute" @do_something
 }
 
 /// TODO: 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 {
+spv.mlir.module Logical GLSL450 {
    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 {
+spv.mlir.module Logical GLSL450 {
    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 {
+spv.mlir.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // expected-error @+1 {{'spv.EntryPoint' invalid execution_model attribute specification: "ContractionOff"}}
    spv.EntryPoint "ContractionOff" @do_nothing
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.ExecutionMode
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
    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 {
+spv.mlir.module Logical GLSL450 {
    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 {
+spv.mlir.module Logical GLSL450 {
    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
 }
 

 // -----
 
 // expected-error @+1 {{expected SSA identifier}}
 spv.func @cannot_have_variadic_arguments(%arg: i32, ...) "None"
 
 // -----
 
 // Nested function
-spv.module Logical GLSL450 {
+spv.mlir.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 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
   spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
 }
 
 // TODO: Fix test case after initialization with normal constant is addressed
-// spv.module Logical GLSL450 {
+// spv.mlir.module Logical GLSL450 {
 //   %0 = spv.Constant 4.0 : f32
 //   // CHECK1: spv.Variable init(%0) : !spv.ptr<f32, Private>
 //   spv.GlobalVariable @var1 init(%0) : !spv.ptr<f32, Private>
 // }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc = 4.0 : f32
   // CHECK: spv.GlobalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
   spv.GlobalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
 }
 
 // -----
 
 // Allow initializers coming from other module-like ops
 spv.SpecConstant @sc = 4.0 : f32
 // CHECK: spv.GlobalVariable @var initializer(@sc)
 spv.GlobalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.GlobalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
   spv.GlobalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
 }
 
 // TODO: Fix test case after initialization with constant is addressed
-// spv.module Logical GLSL450 {
+// spv.mlir.module Logical GLSL450 {
 //   %0 = spv.Constant 4.0 : f32
 //   // CHECK1: spv.GlobalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
 //   spv.GlobalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
 // }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.GlobalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
   spv.GlobalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
   // CHECK: spv.GlobalVariable @var2 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
   spv.GlobalVariable @var2 {built_in = "GlobalInvocationID"} : !spv.ptr<vector<3xi32>, Input>
 }
 
 // -----
 
 // Allow in other module-like ops
 module {
   // CHECK: spv.GlobalVariable
   spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // expected-error @+1 {{expected spv.ptr type}}
   spv.GlobalVariable @var0 : f32
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // expected-error @+1 {{op initializer must be result of a spv.SpecConstant or spv.GlobalVariable op}}
   spv.GlobalVariable @var0 initializer(@var1) : !spv.ptr<f32, Private>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // expected-error @+1 {{storage class cannot be 'Generic'}}
   spv.GlobalVariable @var0 : !spv.ptr<f32, Generic>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // expected-error @+1 {{storage class cannot be 'Function'}}
   spv.GlobalVariable @var0 : !spv.ptr<f32, Function>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   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
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
-// spv.module
+// spv.mlir.module
 //===----------------------------------------------------------------------===//
 
 // Module without capability and extension
-// CHECK: spv.module Logical GLSL450
-spv.module Logical GLSL450 { }
+// CHECK: spv.mlir.module Logical GLSL450
+spv.mlir.module Logical GLSL450 { }
 
 // Module with a name
-// CHECK: spv.module @{{.*}} Logical GLSL450
-spv.module @name Logical GLSL450 { }
+// CHECK: spv.mlir.module @{{.*}} Logical GLSL450
+spv.mlir.module @name Logical GLSL450 { }
 
 // Module with (version, capabilities, extensions) triple
-// CHECK: spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]>
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> { }
+// CHECK: spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]>
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> { }
 
 // Module with additional attributes
-// CHECK: spv.module Logical GLSL450 attributes {foo = "bar"}
-spv.module Logical GLSL450 attributes {foo = "bar"} { }
+// CHECK: spv.mlir.module Logical GLSL450 attributes {foo = "bar"}
+spv.mlir.module Logical GLSL450 attributes {foo = "bar"} { }
 
 // Module with VCE triple and additional attributes
-// CHECK: spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> attributes {foo = "bar"}
-spv.module Logical GLSL450
+// CHECK: spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> attributes {foo = "bar"}
+spv.mlir.module Logical GLSL450
   requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]>
   attributes {foo = "bar"} { }
 
 // Module with explicit spv.mlir.endmodule
-// CHECK: spv.module
-spv.module Logical GLSL450 {
+// CHECK: spv.mlir.module
+spv.mlir.module Logical GLSL450 {
   spv.mlir.endmodule
 }
 
 // Module with function
-// CHECK: spv.module
-spv.module Logical GLSL450 {
+// CHECK: spv.mlir.module
+spv.mlir.module Logical GLSL450 {
   spv.func @do_nothing() -> () "None" {
     spv.Return
   }
 }
 
 // -----
 
 // Missing addressing model
-// expected-error@+1 {{'spv.module' expected valid keyword}}
-spv.module { }
+// expected-error@+1 {{'spv.mlir.module' expected valid keyword}}
+spv.mlir.module { }
 
 // -----
 
 // Wrong addressing model
-// expected-error@+1 {{'spv.module' invalid addressing_model attribute specification: Physical}}
-spv.module Physical { }
+// expected-error@+1 {{'spv.mlir.module' invalid addressing_model attribute specification: Physical}}
+spv.mlir.module Physical { }
 
 // -----
 
 // Missing memory model
-// expected-error@+1 {{'spv.module' expected valid keyword}}
-spv.module Logical { }
+// expected-error@+1 {{'spv.mlir.module' expected valid keyword}}
+spv.mlir.module Logical { }
 
 // -----
 
 // Wrong memory model
-// expected-error@+1 {{'spv.module' invalid memory_model attribute specification: Bla}}
-spv.module Logical Bla { }
+// expected-error@+1 {{'spv.mlir.module' invalid memory_model attribute specification: Bla}}
+spv.mlir.module Logical Bla { }
 
 // -----
 
 // Module with multiple blocks
 // expected-error @+1 {{expects region #0 to have 0 or 1 blocks}}
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
 ^first:
   spv.Return
 ^second:
   spv.Return
 }
 
 // -----
 
 // Module with wrong terminator
 // expected-error@+2 {{expects regions to end with 'spv.mlir.endmodule'}}
 // expected-note@+1 {{in custom textual format, the absence of terminator implies 'spv.mlir.endmodule'}}
-"spv.module"() ({
+"spv.mlir.module"() ({
   %0 = spv.Constant true
 }) {addressing_model = 0 : i32, memory_model = 1 : i32} : () -> ()
 
 // -----
 
 // Use non SPIR-V op inside module
-spv.module Logical GLSL450 {
-  // expected-error @+1 {{'spv.module' can only contain spv.* ops}}
+spv.mlir.module Logical GLSL450 {
+  // expected-error @+1 {{'spv.mlir.module' can only contain spv.* ops}}
   "dialect.op"() : () -> ()
 }
 
 // -----
 
 // Use non SPIR-V op inside function
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @do_nothing() -> () "None" {
-    // expected-error @+1 {{functions in 'spv.module' can only contain spv.* ops}}
+    // expected-error @+1 {{functions in 'spv.mlir.module' can only contain spv.* ops}}
     "dialect.op"() : () -> ()
   }
 }
 
 // -----
 
 // Use external function
-spv.module Logical GLSL450 {
-  // expected-error @+1 {{'spv.module' cannot contain external functions}}
+spv.mlir.module Logical GLSL450 {
+  // expected-error @+1 {{'spv.mlir.module' cannot contain external functions}}
   spv.func @extern() -> () "None"
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.mlir.endmodule
 //===----------------------------------------------------------------------===//
 
 func @module_end_not_in_module() -> () {
   // expected-error @+1 {{op must appear in a module-like op's block}}
   spv.mlir.endmodule
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.mlir.referenceof
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = false
   spv.SpecConstant @sc2 = 42 : i64
   spv.SpecConstant @sc3 = 1.5 : f32
 
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.struct<(i1, i64, f32)>
 
   // CHECK-LABEL: @reference
   spv.func @reference() -> i1 "None" {
func @reference_of_composite() {
   // CHECK: spv.mlir.referenceof @scc : !spv.array<1 x i32>
   %0 = spv.mlir.referenceof @scc : !spv.array<1 x i32>
   %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<1 x i32>
   return
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{expected spv.SpecConstant or spv.SpecConstantComposite symbol}}
     %0 = spv.mlir.referenceof @sc : i32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc = 42 : i32
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}}
     %0 = spv.mlir.referenceof @sc : f32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc = 42 : i32
   spv.SpecConstantComposite @scc (@sc) : !spv.array<1 x i32>
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}}
     %0 = spv.mlir.referenceof @scc : f32
     spv.Return
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.SpecConstant
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.SpecConstant @sc1 = false
   spv.SpecConstant @sc1 = false
   // CHECK: spv.SpecConstant @sc2 spec_id(5) = 42 : i64
   spv.SpecConstant @sc2 spec_id(5) = 42 : i64
   // CHECK: spv.SpecConstant @sc3 = 1.500000e+00 : f32
   spv.SpecConstant @sc3 = 1.5 : f32
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // expected-error @+1 {{SpecId cannot be negative}}
   spv.SpecConstant @sc2 spec_id(-5) = 42 : i64
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // expected-error @+1 {{default value bitwidth disallowed}}
   spv.SpecConstant @sc = 15 : i4
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // expected-error @+1 {{default value can only be a bool, integer, or float scalar}}
   spv.SpecConstant @sc = dense<[2, 3]> : vector<2xi32>
 }
 
 // -----
 
 func @use_in_function() -> () {
   // expected-error @+1 {{op must appear in a module-like op's block}}
   spv.SpecConstant @sc = false
   return
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.SpecConstantComposite
 //===----------------------------------------------------------------------===//
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // expected-error @+1 {{result type must be a composite type}}
   spv.SpecConstantComposite @scc2 (@sc1, @sc2, @sc3) : i32
 }
 
 //===----------------------------------------------------------------------===//
 // spv.SpecConstantComposite (spv.array)
 //===----------------------------------------------------------------------===//
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = 1.5 : f32
   spv.SpecConstant @sc2 = 2.5 : f32
   spv.SpecConstant @sc3 = 3.5 : f32
   // CHECK: spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.array<3 x f32>
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.array<3 x f32>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = false
   spv.SpecConstant @sc2 spec_id(5) = 42 : i64
   spv.SpecConstant @sc3 = 1.5 : f32
   // expected-error @+1 {{has incorrect number of operands: expected 4, but provided 3}}
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.array<4 x f32>
 
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = 1   : i32
   spv.SpecConstant @sc2 = 2.5 : f32
   spv.SpecConstant @sc3 = 3.5 : f32
   // expected-error @+1 {{has incorrect types of operands: expected 'f32', but provided 'i32'}}
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.array<3 x f32>
 }
 
 //===----------------------------------------------------------------------===//
 // spv.SpecConstantComposite (spv.struct)
 //===----------------------------------------------------------------------===//
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = 1   : i32
   spv.SpecConstant @sc2 = 2.5 : f32
   spv.SpecConstant @sc3 = 3.5 : f32
   // CHECK: spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.struct<(i32, f32, f32)>
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.struct<(i32, f32, f32)>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = 1   : i32
   spv.SpecConstant @sc2 = 2.5 : f32
   spv.SpecConstant @sc3 = 3.5 : f32
   // expected-error @+1 {{has incorrect number of operands: expected 2, but provided 3}}
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.struct<(i32, f32)>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = 1.5 : f32
   spv.SpecConstant @sc2 = 2.5 : f32
   spv.SpecConstant @sc3 = 3.5 : f32
   // expected-error @+1 {{has incorrect types of operands: expected 'i32', but provided 'f32'}}
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.struct<(i32, f32, f32)>
 }
 
 //===----------------------------------------------------------------------===//
 // spv.SpecConstantComposite (vector)
 //===----------------------------------------------------------------------===//
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = 1.5 : f32
   spv.SpecConstant @sc2 = 2.5 : f32
   spv.SpecConstant @sc3 = 3.5 : f32
   // CHECK: spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : vector<3xf32>
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : vector<3 x f32>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = false
   spv.SpecConstant @sc2 spec_id(5) = 42 : i64
   spv.SpecConstant @sc3 = 1.5 : f32
   // expected-error @+1 {{has incorrect number of operands: expected 4, but provided 3}}
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : vector<4xf32>
 
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = 1   : i32
   spv.SpecConstant @sc2 = 2.5 : f32
   spv.SpecConstant @sc3 = 3.5 : f32
   // expected-error @+1 {{has incorrect types of operands: expected 'f32', but provided 'i32'}}
   spv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : vector<3xf32>
 }
 
 //===----------------------------------------------------------------------===//
 // spv.SpecConstantComposite (spv.coopmatrix)
 //===----------------------------------------------------------------------===//
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc1 = 1.5 : f32
   // expected-error @+1 {{unsupported composite type}}
   spv.SpecConstantComposite @scc (@sc1) : !spv.coopmatrix<8x16xf32, Device>
 }
 
 //===----------------------------------------------------------------------===//
 // spv.SpecConstantOperation
 //===----------------------------------------------------------------------===//
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo() -> i32 "None" {
     // CHECK: [[LHS:%.*]] = spv.Constant
     %0 = spv.Constant 1: i32
     // CHECK: [[RHS:%.*]] = spv.Constant
     %1 = spv.Constant 1: i32
 
     // CHECK: spv.SpecConstantOperation wraps "spv.IAdd"([[LHS]], [[RHS]]) : (i32, i32) -> i32
     %2 = spv.SpecConstantOperation wraps "spv.IAdd"(%0, %1) : (i32, i32) -> i32
 
     spv.ReturnValue %2 : i32
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @sc = 42 : i32
 
   spv.func @foo() -> i32 "None" {
     // CHECK: [[SC:%.*]] = spv.mlir.referenceof @sc
     %0 = spv.mlir.referenceof @sc : i32
     // CHECK: spv.SpecConstantOperation wraps "spv.ISub"([[SC]], [[SC]]) : (i32, i32) -> i32
     %1 = spv.SpecConstantOperation wraps "spv.ISub"(%0, %0) : (i32, i32) -> i32
     spv.ReturnValue %1 : i32
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo() -> i32 "None" {
     %0 = spv.Constant 1: i32
     // expected-error @+1 {{op expects parent op 'spv.SpecConstantOperation'}}
     spv.mlir.yield %0 : i32
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo() -> () "None" {
     %0 = spv.Variable : !spv.ptr<i32, Function>
 
     // expected-error @+1 {{invalid enclosed op}}
     %1 = spv.SpecConstantOperation wraps "spv.Load"(%0) {memory_access = 0 : i32} : (!spv.ptr<i32, Function>) -> i32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo() -> () "None" {
     %0 = spv.Variable : !spv.ptr<i32, Function>
     %1 = spv.Load "Function" %0 : i32
 
     // expected-error @+1 {{invalid operand, must be defined by a constant operation}}
     %2 = spv.SpecConstantOperation wraps "spv.IAdd"(%1, %1) : (i32, i32) -> i32
 
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/IR/target-env.mlir

func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attributes {
   %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: @module_suitable_extension1
 func @module_suitable_extension1() attributes {
   spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer]>, {}>
 } {
-  // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
+  // CHECK: spv.mlir.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () ->()
   return
 }
 
 // CHECK-LABEL: @module_suitable_extension2
 func @module_suitable_extension2() attributes {
   spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer]>, {}>
 } {
-  // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
+  // CHECK: spv.mlir.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () -> ()
   return
 }
 
 // CHECK-LABEL: @module_missing_extension_mm
 func @module_missing_extension_mm() attributes {
   spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_physical_storage_buffer]>, {}>
 } {
func @module_missing_extension_am() attributes {
   return
 }
 
 // CHECK-LABEL: @module_implied_extension
 func @module_implied_extension() attributes {
   // Version 1.5 implies SPV_KHR_vulkan_memory_model and SPV_KHR_physical_storage_buffer.
   spv.target_env = #spv.target_env<#spv.vce<v1.5, [VulkanMemoryModel, PhysicalStorageBufferAddresses], []>, {}>
 } {
-  // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
+  // CHECK: spv.mlir.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () -> ()
   return
 }

mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir

 // RUN: mlir-opt -test-spirv-module-combiner -split-input-file -verify-diagnostics %s | FileCheck %s
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.SpecConstant @m1_sc
 // CHECK-NEXT:     spv.SpecConstant @m2_sc
 // CHECK-NEXT:     spv.func @variable_init_spec_constant
 // CHECK-NEXT:       spv.mlir.referenceof @m2_sc
 // CHECK-NEXT:       spv.Variable init
 // CHECK-NEXT:       spv.Return
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @m1_sc = 42.42 : f32
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @m2_sc = 42 : i32
   spv.func @variable_init_spec_constant() -> () "None" {
     %0 = spv.mlir.referenceof @m2_sc : i32
     %1 = spv.Variable init(%0) : !spv.ptr<i32, Function>
     spv.Return
   }
 }
 }
 
 // -----
 
 module {
-spv.module Physical64 GLSL450 {
+spv.mlir.module Physical64 GLSL450 {
 }
 
 // expected-error @+1 {{input modules differ in addressing model and/or memory model}}
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
 }
 }
 
 // -----
 
 module {
-spv.module Logical Simple {
+spv.mlir.module Logical Simple {
 }
 
 // expected-error @+1 {{input modules differ in addressing model and/or memory model}}
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
 }
 }

mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict-resolution.mlir

 // RUN: mlir-opt -test-spirv-module-combiner -split-input-file -verify-diagnostics %s | FileCheck %s
 
 // Test basic renaming of conflicting funcOps.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : f32) -> f32 "None" {
     spv.ReturnValue %arg0 : f32
   }
 }
 }
 
 // -----
 
 // Test basic renaming of conflicting funcOps across 3 modules.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo_1
 // CHECK-NEXT:       spv.FAdd
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo_2
 // CHECK-NEXT:       spv.ISub
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : f32) -> f32 "None" {
     %0 = spv.FAdd %arg0, %arg0 : f32
     spv.ReturnValue %0 : f32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     %0 = spv.ISub %arg0, %arg0 : i32
     spv.ReturnValue %0 : i32
   }
 }
 }
 
 // -----
 
 // Test properly updating references to a renamed funcOp.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.FunctionCall @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : f32) -> f32 "None" {
     spv.ReturnValue %arg0 : f32
   }
 
   spv.func @bar(%arg0 : f32) -> f32 "None" {
     %0 = spv.FunctionCall @foo(%arg0) : (f32) ->  (f32)
     spv.ReturnValue %0 : f32
   }
 }
 }
 
 // -----
 
 // Test properly updating references to a renamed funcOp if the functionCallOp
 // preceeds the callee funcOp definition.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.FunctionCall @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @bar(%arg0 : f32) -> f32 "None" {
     %0 = spv.FunctionCall @foo(%arg0) : (f32) ->  (f32)
     spv.ReturnValue %0 : f32
   }
 
   spv.func @foo(%arg0 : f32) -> f32 "None" {
     spv.ReturnValue %arg0 : f32
   }
 }
 }
 
 // -----
 
 // Test properly updating entryPointOp and executionModeOp attached to renamed
 // funcOp.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.EntryPoint "GLCompute" @foo_1
 // CHECK-NEXT:     spv.ExecutionMode @foo_1 "ContractionOff"
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : f32) -> f32 "None" {
     spv.ReturnValue %arg0 : f32
   }
 
   spv.EntryPoint "GLCompute" @foo
   spv.ExecutionMode @foo "ContractionOff"
 }
 }
 
 // -----
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.EntryPoint "GLCompute" @fo
 // CHECK-NEXT:     spv.ExecutionMode @foo "ContractionOff"
 
 // CHECK-NEXT:     spv.func @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.EntryPoint "GLCompute" @foo_1
 // CHECK-NEXT:     spv.ExecutionMode @foo_1 "ContractionOff"
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
   
   spv.EntryPoint "GLCompute" @foo
   spv.ExecutionMode @foo "ContractionOff"
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : f32) -> f32 "None" {
     spv.ReturnValue %arg0 : f32
   }
 
   spv.EntryPoint "GLCompute" @foo
   spv.ExecutionMode @foo "ContractionOff"
 }
 }
 
 // -----
 
 // Resolve conflicting funcOp and globalVariableOp.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.GlobalVariable @foo_1
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 }
 
 // -----
 
 // Resolve conflicting funcOp and globalVariableOp and update the global variable's
 // references.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.GlobalVariable @foo_1
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.mlir.addressof @foo_1
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @bar() -> f32 "None" {
     %0 = spv.mlir.addressof @foo : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOp and funcOp and update the global variable's
 // references.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.GlobalVariable @foo_1
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.mlir.addressof @foo_1
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @bar() -> f32 "None" {
     %0 = spv.mlir.addressof @foo : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting funcOp and specConstantOp.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.SpecConstant @foo_1
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @foo = -5 : i32
 }
 }
 
 // -----
 
 // Resolve conflicting funcOp and specConstantOp and update the spec constant's
 // references.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.SpecConstant @foo_1
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.mlir.referenceof @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @foo = -5 : i32
 
   spv.func @bar() -> i32 "None" {
     %0 = spv.mlir.referenceof @foo : i32 
     spv.ReturnValue %0 : i32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting specConstantOp and funcOp and update the spec constant's
 // references.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.SpecConstant @foo_1
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.mlir.referenceof @foo_1
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @foo = -5 : i32
 
   spv.func @bar() -> i32 "None" {
     %0 = spv.mlir.referenceof @foo : i32
     spv.ReturnValue %0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting funcOp and specConstantCompositeOp.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.SpecConstant @bar
 // CHECK-NEXT:     spv.SpecConstantComposite @foo_1 (@bar, @bar)
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @bar = -5 : i32
   spv.SpecConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 }
 }
 
 // -----
 
 // Resolve conflicting funcOp and specConstantCompositeOp and update the spec
 // constant's references.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.SpecConstant @bar
 // CHECK-NEXT:     spv.SpecConstantComposite @foo_1 (@bar, @bar)
 // CHECK-NEXT:     spv.func @baz
 // CHECK-NEXT:       spv.mlir.referenceof @foo_1
 // CHECK-NEXT:       spv.CompositeExtract
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @bar = -5 : i32
   spv.SpecConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 
   spv.func @baz() -> i32 "None" {
     %0 = spv.mlir.referenceof @foo : !spv.array<2 x i32>
     %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<2 x i32>
     spv.ReturnValue %1 : i32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting specConstantCompositeOp and funcOp and update the spec
 // constant's references.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.SpecConstant @bar
 // CHECK-NEXT:     spv.SpecConstantComposite @foo_1 (@bar, @bar)
 // CHECK-NEXT:     spv.func @baz
 // CHECK-NEXT:       spv.mlir.referenceof @foo_1
 // CHECK-NEXT:       spv.CompositeExtract
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @bar = -5 : i32
   spv.SpecConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 
   spv.func @baz() -> i32 "None" {
     %0 = spv.mlir.referenceof @foo : !spv.array<2 x i32>
     %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<2 x i32>
     spv.ReturnValue %1 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting spec constants and funcOps and update the spec constant's
 // references.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.SpecConstant @bar_1
 // CHECK-NEXT:     spv.SpecConstantComposite @foo_2 (@bar_1, @bar_1)
 // CHECK-NEXT:     spv.func @baz
 // CHECK-NEXT:       spv.mlir.referenceof @foo_2
 // CHECK-NEXT:       spv.CompositeExtract
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @bar = -5 : i32
   spv.SpecConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 
   spv.func @baz() -> i32 "None" {
     %0 = spv.mlir.referenceof @foo : !spv.array<2 x i32>
     %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<2 x i32>
     spv.ReturnValue %1 : i32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 
   spv.func @bar(%arg0 : f32) -> f32 "None" {
     spv.ReturnValue %arg0 : f32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOps.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.GlobalVariable @foo_1 bind(1, 0)
 
 // CHECK-NEXT:     spv.GlobalVariable @foo bind(2, 0)
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(2, 0) : !spv.ptr<f32, Input>
 }
 }
 
 // -----
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.GlobalVariable @foo_1 built_in("GlobalInvocationId")
 
 // CHECK-NEXT:     spv.GlobalVariable @foo built_in("LocalInvocationId")
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOp and specConstantOp.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.GlobalVariable @foo_1
 
 // CHECK-NEXT:     spv.SpecConstant @foo
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @foo = -5 : i32
 }
 }
 
 // -----
 
 // Resolve conflicting specConstantOp and globalVariableOp.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.SpecConstant @foo_1
 
 // CHECK-NEXT:     spv.GlobalVariable @foo
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @foo = -5 : i32
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOp and specConstantCompositeOp.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.GlobalVariable @foo_1
 
 // CHECK-NEXT:     spv.SpecConstant @bar
 // CHECK-NEXT:     spv.SpecConstantComposite @foo (@bar, @bar)
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @bar = -5 : i32
   spv.SpecConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOp and specConstantComposite.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.SpecConstant @bar
 // CHECK-NEXT:     spv.SpecConstantComposite @foo_1 (@bar, @bar)
 
 // CHECK-NEXT:     spv.GlobalVariable @foo
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @bar = -5 : i32
   spv.SpecConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 }

mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir

 // RUN: mlir-opt -test-spirv-module-combiner -split-input-file -verify-diagnostics %s | FileCheck %s
 
 // Deduplicate 2 global variables with the same descriptor set and binding.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.GlobalVariable @foo
 
 // CHECK-NEXT:     spv.func @use_foo
 // CHECK-NEXT:       spv.mlir.addressof @foo
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @use_bar
 // CHECK-NEXT:       spv.mlir.addressof @foo
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.FAdd
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @use_foo() -> f32 "None" {
     %0 = spv.mlir.addressof @foo : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @bar bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @use_bar() -> f32 "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     %2 = spv.FAdd %1, %1 : f32
     spv.ReturnValue %2 : f32
   }
 }
 }
 
 // -----
 
 // Deduplicate 2 global variables with the same descriptor set and binding but different types.
 
 // CHECK:      module {
-// CHECK-NEXT: spv.module Logical GLSL450 {
+// CHECK-NEXT: spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:   spv.GlobalVariable @foo bind(1, 0)
 
 // CHECK-NEXT:   spv.GlobalVariable @bar bind(1, 0)
 
 // CHECK-NEXT:   spv.func @use_bar
 // CHECK-NEXT:     spv.mlir.addressof @bar
 // CHECK-NEXT:     spv.Load
 // CHECK-NEXT:     spv.ReturnValue
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<i32, Input>
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @bar bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @use_bar() -> f32 "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 }
 
 // -----
 
 // Deduplicate 2 global variables with the same built-in attribute.
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.GlobalVariable @foo built_in("GlobalInvocationId")
 // CHECK-NEXT:     spv.func @use_bar
 // CHECK-NEXT:       spv.mlir.addressof @foo
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @foo built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @bar built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 
   spv.func @use_bar() -> vector<3xi32> "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<vector<3xi32>, Input>
     %1 = spv.Load "Input" %0 : vector<3xi32>
     spv.ReturnValue %1 : vector<3xi32>
   }
 }
 }
 
 // -----
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.SpecConstant @foo spec_id(5)
 
 // CHECK-NEXT:     spv.func @use_foo()
 // CHECK-NEXT:       %0 = spv.mlir.referenceof @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @use_bar()
 // CHECK-NEXT:       %0 = spv.mlir.referenceof @foo
 // CHECK-NEXT:       spv.FAdd
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @foo spec_id(5) = 1. : f32
 
   spv.func @use_foo() -> (f32) "None" {
     %0 = spv.mlir.referenceof @foo : f32
     spv.ReturnValue %0 : f32
   }
 }
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @bar spec_id(5) = 1. : f32
 
   spv.func @use_bar() -> (f32) "None" {
     %0 = spv.mlir.referenceof @bar : f32
     %1 = spv.FAdd %0, %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 }
 
 // -----
 
 // CHECK:      module {
-// CHECK-NEXT:   spv.module Logical GLSL450 {
+// CHECK-NEXT:   spv.mlir.module Logical GLSL450 {
 // CHECK-NEXT:     spv.SpecConstant @bar spec_id(5)
 
 // CHECK-NEXT:     spv.func @foo(%arg0: f32)
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo_different_body(%arg0: f32)
 // CHECK-NEXT:       spv.mlir.referenceof

 
 // CHECK-NEXT:     spv.func @kernel_different_attr
 // CHECK-NEXT:       spv.Return
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT:   }
 
 module {
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.SpecConstant @bar spec_id(5) = 1. : f32
 
   spv.func @foo(%arg0: f32) -> (f32) "None" {
     spv.ReturnValue %arg0 : f32
   }
 
   spv.func @foo_duplicate(%arg0: f32) -> (f32) "None" {
     spv.ReturnValue %arg0 : f32

mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, {}>
 } {
-  spv.module Physical64 OpenCL {
-    // CHECK-LABEL: spv.module
+  spv.mlir.module Physical64 OpenCL {
+    // CHECK-LABEL: spv.mlir.module
     //       CHECK:   spv.func [[FN:@.*]]({{%.*}}: f32, {{%.*}}: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>
     //       CHECK:   spv.EntryPoint "Kernel" [[FN]]
     //       CHECK:   spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
     spv.func @kernel(
       %arg0: f32,
       %arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>) "None"
     attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
       spv.Return
     }
   }
 }

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

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
 } {
 
-// CHECK-LABEL: spv.module
-spv.module Logical GLSL450 {
+// CHECK-LABEL: spv.mlir.module
+spv.mlir.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, stride=4> [0])>, StorageBuffer>
   // CHECK:    spv.func [[FN:@.*]]()
   spv.func @kernel(
     %arg0: f32
            {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>},
     %arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, StorageBuffer>
            {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}) "None"
   attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     // CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]]
     // CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[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
-} // end spv.module
+} // end spv.mlir.module
 
 } // end module

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

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
 } {
 
-// CHECK-LABEL: spv.module
-spv.module Logical GLSL450 {
+// CHECK-LABEL: spv.mlir.module
+spv.mlir.module Logical GLSL450 {
   // CHECK-DAG: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
   spv.GlobalVariable @__builtin_var_WorkgroupSize__ built_in("WorkgroupSize") : !spv.ptr<vector<3xi32>, Input>
   // CHECK-DAG: spv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
   spv.GlobalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
   // CHECK-DAG: spv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
   spv.GlobalVariable @__builtin_var_LocalInvocationId__ built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
   // CHECK-DAG: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   spv.GlobalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
attributes  {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     %42 = spv.FAdd %39, %41 : f32
     // CHECK: spv.AccessChain [[ARG2]]
     %43 = spv.AccessChain %arg2[%c0, %36, %37] : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32>>)>, StorageBuffer>, i32, i32, i32
     spv.Store "StorageBuffer" %43, %42 : f32
     spv.Return
   }
   // CHECK: spv.EntryPoint "GLCompute" [[FN]], [[WORKGROUPID]], [[LOCALINVOCATIONID]], [[NUMWORKGROUPS]], [[WORKGROUPSIZE]]
   // CHECK-NEXT: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
-} // end spv.module
+} // end spv.mlir.module
 
 } // end module

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

-// RUN: mlir-opt %s -split-input-file -pass-pipeline='spv.module(inline{default-pipeline=''})' | FileCheck %s
+// RUN: mlir-opt %s -split-input-file -pass-pipeline='spv.mlir.module(inline{default-pipeline=''})' | FileCheck %s
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @callee() "None" {
     spv.Return
   }
 
   // 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 {
+spv.mlir.module Logical GLSL450 {
   spv.func @callee() -> i32 "None" {
     %0 = spv.Constant 42 : i32
     spv.ReturnValue %0 : 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.mlir.module Logical GLSL450 {
   spv.GlobalVariable @data bind(0, 0) : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
   spv.func @callee() "None" {
     %0 = spv.mlir.addressof @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>, i32, i32
     spv.Branch ^next
 
   ^next:
spv.func @calling_multi_block_ret_func() "None" {
     spv.Return
   }
 }
 
 // TODO: calling_multi_block_retval_func
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @callee(%cond : i1) -> () "None" {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
       spv.Return
     ^merge:
       spv.mlir.merge
     }
     spv.Return
   }
 
   // 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 {
+spv.mlir.module Logical GLSL450 {
   spv.func @callee(%cond : i1) -> () "None" {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
       spv.Branch ^merge
     ^merge:
       spv.mlir.merge
     }
spv.func @calling_selection_no_ret_func() "None" {
     // CHECK-NEXT:   spv.mlir.merge
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @callee(%cond : i1) -> () "None" {
     spv.loop {
       spv.Branch ^header
     ^header:
       spv.BranchConditional %cond, ^body, ^merge
     ^body:
       spv.Return
     ^continue:
spv.func @calling_loop_ret_func() "None" {
     // CHECK: spv.FunctionCall
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @callee(%cond : i1) -> () "None" {
     spv.loop {
       spv.Branch ^header
     ^header:
       spv.BranchConditional %cond, ^body, ^merge
     ^body:
       spv.Branch ^continue
     ^continue:
spv.func @calling_loop_no_ret_func() "None" {
     // CHECK-NEXT:   spv.mlir.merge
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.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: @inline_into_selection_region
   spv.func @inline_into_selection_region() "None" {
     %1 = spv.Constant 0 : i32
     // CHECK-DAG: [[ADDRESS_ARG0:%.*]] = spv.mlir.addressof @arg_0
     // CHECK-DAG: [[ADDRESS_ARG1:%.*]] = spv.mlir.addressof @arg_1

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

 // RUN: mlir-opt -decorate-spirv-composite-type-layout -split-input-file -verify-diagnostics %s -o - | FileCheck %s
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
   spv.GlobalVariable @var0 bind(0,1) : !spv.ptr<!spv.struct<(i32, !spv.struct<(f32, i32)>, f32)>, Uniform>
 
   // CHECK: spv.GlobalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<(!spv.array<64 x i32, stride=4> [0], f32 [256])>, StorageBuffer>
   spv.GlobalVariable @var1 bind(0,2) : !spv.ptr<!spv.struct<(!spv.array<64xi32>, f32)>, StorageBuffer>
 
   // CHECK: spv.GlobalVariable @var2 bind(1, 0) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64 x i32, stride=4> [0], f32 [256])> [0], i32 [260])>, StorageBuffer>
   spv.GlobalVariable @var2 bind(1,0) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64xi32>, f32)>, i32)>, StorageBuffer>
spv.func @kernel() -> () "None" {
     // 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>, i32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(i1 [0], i1 [1], f64 [8])> [0], i1 [16])> [0], i1 [24])> [0], i1 [32])> [0], i1 [40])>, Uniform>
   spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(i1, i1, f64)>, i1)>, i1)>, i1)>, i1)>, Uniform>
 
   // CHECK: spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(!spv.struct<(i16 [0], !spv.struct<(i1 [0], f64 [8])> [8], f32 [24])> [0], f32 [32])>, Uniform>
   spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(!spv.struct<(i16, !spv.struct<(i1, f64)>, f32)>, f32)>, Uniform>
 
   // CHECK: spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(!spv.struct<(i16 [0], !spv.struct<(i1 [0], !spv.array<16 x !spv.array<16 x i64, stride=8>, stride=128> [8])> [8], f32 [2064])> [0], f32 [2072])>, Uniform>
   spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(!spv.struct<(i16, !spv.struct<(i1, !spv.array<16x!spv.array<16xi64>>)>, f32)>, f32)>, Uniform>
spv.mlir.module Logical GLSL450 {
   spv.GlobalVariable @var6 : !spv.ptr<!spv.struct<(i1, !spv.struct<(i64, i32, i16, i8, i1)>, i1)>, Uniform>
 
   // CHECK: spv.GlobalVariable @var7 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(!spv.struct<(i1 [0], i64 [8])> [0], i1 [16])> [8], i1 [32])>, Uniform>
   spv.GlobalVariable @var7 : !spv.ptr<!spv.struct<(i1, !spv.struct<(!spv.struct<(i1, i64)>, i1)>, i1)>, Uniform>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(vector<2xi32> [0], f32 [8])>, StorageBuffer>
   spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(vector<2xi32>, f32)>, StorageBuffer>
 
   // CHECK: spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(vector<3xi32> [0], f32 [12])>, StorageBuffer>
   spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(vector<3xi32>, f32)>, StorageBuffer>
 
   // CHECK: spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(vector<4xi32> [0], f32 [16])>, StorageBuffer>
   spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(vector<4xi32>, f32)>, StorageBuffer>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.GlobalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<(!spv.struct<()> [0])>, StorageBuffer>
   spv.GlobalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<(!spv.struct<()>)>, StorageBuffer>
 
   // CHECK: spv.GlobalVariable @arrayType : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, StorageBuffer>
   spv.GlobalVariable @arrayType : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, StorageBuffer>
 
   // CHECK: spv.GlobalVariable @InputStorage : !spv.ptr<!spv.struct<(!spv.array<256 x f32>)>, Input>
   spv.GlobalVariable @InputStorage : !spv.ptr<!spv.struct<(!spv.array<256xf32>)>, Input>
 
   // CHECK: spv.GlobalVariable @customLayout : !spv.ptr<!spv.struct<(f32 [256], i32 [512])>, Uniform>
   spv.GlobalVariable @customLayout : !spv.ptr<!spv.struct<(f32 [256], i32 [512])>, Uniform>
 
   // CHECK:  spv.GlobalVariable @emptyStruct : !spv.ptr<!spv.struct<()>, Uniform>
   spv.GlobalVariable @emptyStruct : !spv.ptr<!spv.struct<()>, Uniform>
 }
 
 // -----
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   // CHECK: spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(i32 [0])>, PushConstant>
   spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(i32)>, PushConstant>
   // CHECK: spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(i32 [0])>, PhysicalStorageBuffer>
   spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(i32)>, PhysicalStorageBuffer>
 }

mlir/test/Dialect/SPIRV/Transforms/rewrite-inserts.mlir

 // RUN: mlir-opt -spirv-rewrite-inserts -split-input-file -verify-diagnostics %s -o - | FileCheck %s
 
-spv.module Logical GLSL450 {
+spv.mlir.module Logical GLSL450 {
   spv.func @rewrite(%value0 : f32, %value1 : f32, %value2 : f32, %value3 : i32, %value4: !spv.array<3xf32>) -> vector<3xf32> "None" {
     %0 = spv.undef : vector<3xf32>
     // CHECK: spv.CompositeConstruct {{%.*}}, {{%.*}}, {{%.*}} : vector<3xf32>
     %1 = spv.CompositeInsert %value0, %0[0 : i32] : f32 into vector<3xf32>
     %2 = spv.CompositeInsert %value1, %1[1 : i32] : f32 into vector<3xf32>
     %3 = spv.CompositeInsert %value2, %2[2 : i32] : f32 into vector<3xf32>
 
     %4 = spv.undef : !spv.array<4xf32>

mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir

 // RUN: mlir-opt -spirv-update-vce %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // Version
 //===----------------------------------------------------------------------===//
 
 // Test deducing minimal version.
 // spv.IAdd is available from v1.0.
 
 // CHECK: requires #spv.vce<v1.0, [Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.5, [Shader], []>, {}>
 } {
   spv.func @iadd(%val : i32) -> i32 "None" {
     %0 = spv.IAdd %val, %val: i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test deducing minimal version.
 // spv.GroupNonUniformBallot is available since v1.3.
 
 // CHECK: requires #spv.vce<v1.3, [GroupNonUniformBallot, Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.5, [Shader, GroupNonUniformBallot], []>, {}>
 } {
   spv.func @group_non_uniform_ballot(%predicate : i1) -> vector<4xi32> "None" {
     %0 = spv.GroupNonUniformBallot Workgroup %predicate : vector<4xi32>
     spv.ReturnValue %0: vector<4xi32>
   }
 }
 
 //===----------------------------------------------------------------------===//
 // Capability
 //===----------------------------------------------------------------------===//
 
 // Test minimal capabilities.
 
 // CHECK: requires #spv.vce<v1.0, [Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader, Float16, Float64, Int16, Int64, VariablePointers], []>, {}>
 } {
   spv.func @iadd(%val : i32) -> i32 "None" {
     %0 = spv.IAdd %val, %val: i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test deducing implied capability.
 // AtomicStorage implies Shader.
 
 // CHECK: requires #spv.vce<v1.0, [Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [AtomicStorage], []>, {}>
 } {
   spv.func @iadd(%val : i32) -> i32 "None" {
     %0 = spv.IAdd %val, %val: i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test selecting the capability available in the target environment.
 // spv.GroupNonUniform op itself can be enabled via any of
 // * GroupNonUniformArithmetic
 // * GroupNonUniformClustered
 // * GroupNonUniformPartitionedNV
 // Its 'Reduce' group operation can be enabled via any of
 // * Kernel
 // * GroupNonUniformArithmetic
 // * GroupNonUniformBallot
 
 // CHECK: requires #spv.vce<v1.3, [GroupNonUniformArithmetic, Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
 } {
   spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
     %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 }
 
 // CHECK: requires #spv.vce<v1.3, [GroupNonUniformClustered, GroupNonUniformBallot, Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.3, [Shader, GroupNonUniformClustered, GroupNonUniformBallot], []>, {}>
 } {
   spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
     %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test type required capabilities
 
 // Using 8-bit integers in non-interface storage class requires Int8.
 // CHECK: requires #spv.vce<v1.0, [Int8, Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.3, [Shader, Int8], []>, {}>
 } {
   spv.func @iadd_function(%val : i8) -> i8 "None" {
     %0 = spv.IAdd %val, %val : i8
     spv.ReturnValue %0: i8
   }
 }
 
 // Using 16-bit floats in non-interface storage class requires Float16.
 // CHECK: requires #spv.vce<v1.0, [Float16, Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.3, [Shader, Float16], []>, {}>
 } {
   spv.func @fadd_function(%val : f16) -> f16 "None" {
     %0 = spv.FAdd %val, %val : f16
     spv.ReturnValue %0: f16
   }
 }
 
 // Using 16-element vectors requires Vector16.
 // CHECK: requires #spv.vce<v1.0, [Vector16, Shader], []>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.3, [Shader, Vector16], []>, {}>
 } {
   spv.func @iadd_v16_function(%val : vector<16xi32>) -> vector<16xi32> "None" {
     %0 = spv.IAdd %val, %val : vector<16xi32>
     spv.ReturnValue %0: vector<16xi32>
   }
 }
 
 //===----------------------------------------------------------------------===//
 // Extension
 //===----------------------------------------------------------------------===//
 
 // Test deducing minimal extensions.
 // spv.SubgroupBallotKHR requires the SPV_KHR_shader_ballot extension.
 
 // CHECK: requires #spv.vce<v1.0, [SubgroupBallotKHR, Shader], [SPV_KHR_shader_ballot]>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader, SubgroupBallotKHR],
              [SPV_KHR_shader_ballot, SPV_KHR_shader_clock, SPV_KHR_variable_pointers]>, {}>
 } {
   spv.func @subgroup_ballot(%predicate : i1) -> vector<4xi32> "None" {
     %0 = spv.SubgroupBallotKHR %predicate: vector<4xi32>
     spv.ReturnValue %0: vector<4xi32>
   }
 }
 
 // Test deducing implied extension.
 // Vulkan memory model requires SPV_KHR_vulkan_memory_model, which is enabled
 // implicitly by v1.5.
 
 // CHECK: requires #spv.vce<v1.0, [VulkanMemoryModel], [SPV_KHR_vulkan_memory_model]>
-spv.module Logical Vulkan attributes {
+spv.mlir.module Logical Vulkan attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.5, [Shader, VulkanMemoryModel], []>, {}>
 } {
   spv.func @iadd(%val : i32) -> i32 "None" {
     %0 = spv.IAdd %val, %val: i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test type required extensions
 
 // Using 8-bit integers in interface storage class requires additional
 // extensions and capabilities.
 // CHECK: requires #spv.vce<v1.0, [StorageBuffer16BitAccess, Shader, Int16], [SPV_KHR_16bit_storage, SPV_KHR_storage_buffer_storage_class]>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.3, [Shader, StorageBuffer16BitAccess, Int16], []>, {}>
 } {
   spv.func @iadd_storage_buffer(%ptr : !spv.ptr<i16, StorageBuffer>) -> i16 "None" {
     %0 = spv.Load "StorageBuffer" %ptr : i16
     %1 = spv.IAdd %0, %0 : i16
     spv.ReturnValue %1: i16
   }
 }
 
 // Complicated nested types
 // * Buffer requires ImageBuffer or SampledBuffer.
 // * Rg32f requires StorageImageExtendedFormats.
 // CHECK: requires #spv.vce<v1.0, [UniformAndStorageBuffer8BitAccess, StorageUniform16, Shader, ImageBuffer, StorageImageExtendedFormats], [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>
-spv.module Logical GLSL450 attributes {
+spv.mlir.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.5, [Shader, UniformAndStorageBuffer8BitAccess, StorageBuffer16BitAccess, StorageUniform16, Int16, ImageBuffer, StorageImageExtendedFormats], []>,
     {}>
 } {
   spv.GlobalVariable @data : !spv.ptr<!spv.struct<(i8 [0], f16 [2], i64 [4])>, Uniform>
   spv.GlobalVariable @img  : !spv.ptr<!spv.image<f32, Buffer, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Rg32f>, UniformConstant>
 }

mlir/test/Target/SPIRV/arithmetic-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @fmul(%arg0 : f32, %arg1 : f32) "None" {
     // CHECK: {{%.*}}= spv.FMul {{%.*}}, {{%.*}} : f32
     %0 = spv.FMul %arg0, %arg1 : f32
     spv.Return
   }
   spv.func @fadd(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: {{%.*}} = spv.FAdd {{%.*}}, {{%.*}} : vector<4xf32>
     %0 = spv.FAdd %arg0, %arg1 : vector<4xf32>

mlir/test/Target/SPIRV/array.mlir

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

mlir/test/Target/SPIRV/atomic-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @atomic_compare_exchange_weak
   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>

mlir/test/Target/SPIRV/barrier-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @memory_barrier_0() -> () "None" {
     // CHECK: spv.MemoryBarrier Device, "Release|UniformMemory"
     spv.MemoryBarrier Device, "Release|UniformMemory"
     spv.Return
   }
   spv.func @memory_barrier_1() -> () "None" {
     // CHECK: spv.MemoryBarrier Subgroup, "AcquireRelease|SubgroupMemory"
     spv.MemoryBarrier Subgroup, "AcquireRelease|SubgroupMemory"

mlir/test/Target/SPIRV/bit-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @bitcount(%arg: i32) -> i32 "None" {
     // CHECK: spv.BitCount {{%.*}} : i32
     %0 = spv.BitCount %arg : i32
     spv.ReturnValue %0 : i32
   }
   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

mlir/test/Target/SPIRV/cast-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @bit_cast(%arg0 : f32) "None" {
     // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : f32 to i32
     %0 = spv.Bitcast %arg0 : f32 to i32
     // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : i32 to si32
     %1 = spv.Bitcast %0 : i32 to si32
     // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : si32 to i32
     %2 = spv.Bitcast %1 : si32 to ui32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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
   }
   spv.func @convert_f64_to_s32(%arg0 : f64) -> i32 "None" {
     // CHECK: {{%.*}} = spv.ConvertFToS {{%.*}} : f64 to i32
     %0 = spv.ConvertFToS %arg0 : f64 to i32

mlir/test/Target/SPIRV/composite-op.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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)>)>
   }
   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>

mlir/test/Target/SPIRV/constant.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @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>

mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [CooperativeMatrixNV], [SPV_NV_cooperative_matrix]> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [CooperativeMatrixNV], [SPV_NV_cooperative_matrix]> {
   // CHECK-LABEL: @cooperative_matrix_load
   spv.func @cooperative_matrix_load(%ptr : !spv.ptr<i32, StorageBuffer>, %stride : i32, %b : i1) "None" {
     // CHECK: {{%.*}} = spv.CooperativeMatrixLoadNV {{%.*}}, {{%.*}}, {{%.*}} : !spv.ptr<i32, StorageBuffer> as !spv.coopmatrix<16x8xi32, Workgroup>
     %0 = spv.CooperativeMatrixLoadNV %ptr, %stride, %b : !spv.ptr<i32, StorageBuffer> as !spv.coopmatrix<16x8xi32, Workgroup>
     spv.Return
   }
 
   // CHECK-LABEL: @cooperative_matrix_load_memaccess

mlir/test/Target/SPIRV/debug.mlir

 // RUN: mlir-translate -test-spirv-roundtrip-debug -mlir-print-debuginfo -mlir-print-local-scope %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: loc({{".*debug.mlir"}}:5:3)
   spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<f32, Input>
   spv.func @arithmetic(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: loc({{".*debug.mlir"}}:8:10)
     %0 = spv.FAdd %arg0, %arg1 : vector<4xf32>
     // CHECK: loc({{".*debug.mlir"}}:10:10)
     %1 = spv.FNegate %arg0 : vector<4xf32>
     spv.Return

mlir/test/Target/SPIRV/decorations.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: location = 0 : i32
   spv.GlobalVariable @var {location = 0 : i32} : !spv.ptr<vector<4xf32>, Input>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: no_perspective
   spv.GlobalVariable @var {no_perspective} : !spv.ptr<vector<4xf32>, Input>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: flat
   spv.GlobalVariable @var {flat} : !spv.ptr<si32, Input>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: aliased
   // CHECK: aliased
   spv.GlobalVariable @var1 bind(0, 0) {aliased} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
   spv.GlobalVariable @var2 bind(0, 0) {aliased} : !spv.ptr<!spv.struct<(vector<4xf32>[0])>, StorageBuffer>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: non_readable
   spv.GlobalVariable @var bind(0, 0) {non_readable} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: non_writable
   spv.GlobalVariable @var bind(0, 0) {non_writable} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: restrict
   spv.GlobalVariable @var bind(0, 0) {restrict} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
 }
 

mlir/test/Target/SPIRV/entry-point.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK:       spv.GlobalVariable @var2 : !spv.ptr<f32, Input>
   // CHECK-NEXT:  spv.GlobalVariable @var3 : !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>
   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/Target/SPIRV/execution-mode.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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/Target/SPIRV/function-call.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.GlobalVariable @var1 : !spv.ptr<!spv.array<4xf32>, Input>
   spv.func @fmain() -> i32 "None" {
     %0 = spv.Constant 16 : i32
     %1 = spv.mlir.addressof @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>) ->  ()

mlir/test/Target/SPIRV/global-variable.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 // CHECK:      spv.GlobalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
 // CHECK-NEXT: spv.GlobalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
 // CHECK-NEXT: spv.GlobalVariable @var2 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 // CHECK-NEXT: spv.GlobalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.GlobalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
   spv.GlobalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
   spv.GlobalVariable @var2 {built_in = "GlobalInvocationId"} : !spv.ptr<vector<3xi32>, Input>
   spv.GlobalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK:         spv.GlobalVariable @var1 : !spv.ptr<f32, Input>
   // CHECK-NEXT:    spv.GlobalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
   spv.GlobalVariable @var1 : !spv.ptr<f32, Input>
   spv.GlobalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.GlobalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
   spv.func @foo() "None" {
     // CHECK: %[[ADDR:.*]] = spv.mlir.addressof @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
     %0 = spv.mlir.addressof @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>, i32
     spv.Return
   }
 }

mlir/test/Target/SPIRV/glsl-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @fmul(%arg0 : f32, %arg1 : f32, %arg2 : i32) "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

mlir/test/Target/SPIRV/group-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @subgroup_ballot
   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>
   }
   // CHECK-LABEL: @group_broadcast_1
   spv.func @group_broadcast_1(%value: f32, %localid: i32 ) -> f32 "None" {

mlir/test/Target/SPIRV/image.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: !spv.ptr<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>, UniformConstant>
   spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>, UniformConstant>
 
   // CHECK: !spv.ptr<!spv.image<si32, Cube, IsDepth, NonArrayed, SingleSampled, NeedSampler, R8ui>, UniformConstant>
   spv.GlobalVariable @var1 : !spv.ptr<!spv.image<si32, Cube, IsDepth, NonArrayed, SingleSampled, NeedSampler, R8ui>, UniformConstant>
 
   // CHECK: !spv.ptr<!spv.image<i32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>, UniformConstant>
   spv.GlobalVariable @var2 : !spv.ptr<!spv.image<i32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>, UniformConstant>
 }

mlir/test/Target/SPIRV/logical-ops.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @iequal_scalar(%arg0: i32, %arg1: i32)  "None" {
     // CHECK: {{.*}} = spv.IEqual {{.*}}, {{.*}} : i32
     %0 = spv.IEqual %arg0, %arg1 : i32
     spv.Return
   }
   spv.func @inotequal_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.INotEqual {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.INotEqual %arg0, %arg1 : vector<4xi32>
spv.func @cmpf(%arg0 : f32, %arg1 : f32) "None" {
     // CHCK: spv.IsInf
     %16 = spv.IsInf %arg1 : f32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.SpecConstant @condition_scalar = true
   spv.func @select() -> () "None" {
     %0 = spv.Constant 4.0 : f32
     %1 = spv.Constant 5.0 : f32
     %2 = spv.mlir.referenceof @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>

mlir/test/Target/SPIRV/loop.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 // Single loop
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // for (int i = 0; i < count; ++i) {}
   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.mlir.merge
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.GlobalVariable @GV1 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
   spv.GlobalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
   spv.func @loop_kernel() "None" {
     %0 = spv.mlir.addressof @GV1 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
     %1 = spv.Constant 0 : i32
     %2 = spv.AccessChain %0[%1] : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>, i32
     %3 = spv.mlir.addressof @GV2 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
     %5 = spv.AccessChain %3[%1] : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>, i32
// CHECK-NEXT:     spv.mlir.merge
   spv.EntryPoint "GLCompute" @loop_kernel
   spv.ExecutionMode @loop_kernel "LocalSize", 1, 1, 1
 }
 
 // -----
 
 // Nested loop
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // for (int i = 0; i < count; ++i) {
   //   for (int j = 0; j < count; ++j) { }
   // }
   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>

mlir/test/Target/SPIRV/matrix.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @matrix_access_chain
   spv.func @matrix_access_chain(%arg0 : !spv.ptr<!spv.matrix<3 x vector<3xf32>>, Function>, %arg1 : i32) -> !spv.ptr<vector<3xf32>, Function> "None" {
     // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.matrix<3 x vector<3xf32>>, Function>
     %0 = spv.AccessChain %arg0[%arg1] : !spv.ptr<!spv.matrix<3 x vector<3xf32>>,Function>, i32
     spv.ReturnValue %0 : !spv.ptr<vector<3xf32>, Function>
   }
 
   // CHECK-LABEL: @matrix_times_scalar_1
spv.func @matrix_times_matrix_2(%arg0: !spv.matrix<3 x vector<2xf32>>, %arg1: !spv.matrix<2 x vector<3xf32>>) -> !spv.matrix<2 x vector<2xf32>> "None"{
     // CHECK: {{%.*}} = spv.MatrixTimesMatrix {{%.*}}, {{%.*}} : !spv.matrix<3 x vector<2xf32>>, !spv.matrix<2 x vector<3xf32>> -> !spv.matrix<2 x vector<2xf32>>
     %result = spv.MatrixTimesMatrix %arg0, %arg1 : !spv.matrix<3 x vector<2xf32>>, !spv.matrix<2 x vector<3xf32>> -> !spv.matrix<2 x vector<2xf32>>
     spv.ReturnValue %result : !spv.matrix<2 x vector<2xf32>>
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr<!spv.matrix<3 x vector<3xf32>>, StorageBuffer>
   spv.GlobalVariable @var0 : !spv.ptr<!spv.matrix<3 x vector<3xf32>>, StorageBuffer>
 
   // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr<!spv.matrix<2 x vector<3xf32>>, StorageBuffer>
   spv.GlobalVariable @var1 : !spv.ptr<!spv.matrix<2 x vector<3xf32>>, StorageBuffer>
 
   // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr<!spv.matrix<4 x vector<4xf16>>, StorageBuffer>
   spv.GlobalVariable @var2 : !spv.ptr<!spv.matrix<4 x vector<4xf16>>, StorageBuffer>
 }

mlir/test/Target/SPIRV/memory-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 // 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 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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>, i32
     %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, i32, i32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @load_store_zero_rank_float(%arg0: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>) "None" {
     // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>
     // CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : f32
     %0 = spv.Constant 0 : i32
     %1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
     %2 = spv.Load "StorageBuffer" %1 : f32
 
     // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>
spv.func @load_store_zero_rank_int(%arg0: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>) "None" {
     %4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
     spv.Store "StorageBuffer" %4, %2 : i32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @copy_memory_simple() "None" {
     %0 = spv.Variable : !spv.ptr<f32, Function>
     %1 = spv.Variable : !spv.ptr<f32, Function>
     // CHECK: spv.CopyMemory "Function" %{{.*}}, "Function" %{{.*}} : f32
     spv.CopyMemory "Function" %0, "Function" %1 : f32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @copy_memory_different_storage_classes(%in : !spv.ptr<!spv.array<4xf32>, Input>, %out : !spv.ptr<!spv.array<4xf32>, Output>) "None" {
     // CHECK: spv.CopyMemory "Output" %{{.*}}, "Input" %{{.*}} : !spv.array<4 x f32>
     spv.CopyMemory "Output" %out, "Input" %in : !spv.array<4xf32>
     spv.Return
   }
 }
 
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @copy_memory_with_access_operands() "None" {
     %0 = spv.Variable : !spv.ptr<f32, Function>
     %1 = spv.Variable : !spv.ptr<f32, Function>
     // CHECK: spv.CopyMemory "Function" %{{.*}}, "Function" %{{.*}} ["Aligned", 4] : f32
     spv.CopyMemory "Function" %0, "Function" %1 ["Aligned", 4] : f32
 
     // CHECK: spv.CopyMemory "Function" %{{.*}}, "Function" %{{.*}} ["Volatile"] : f32
     spv.CopyMemory "Function" %0, "Function" %1 ["Volatile"] : f32

mlir/test/Target/SPIRV/module.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-// CHECK:      spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+// CHECK:      spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 // CHECK-NEXT:   spv.func @foo() "Inline" {
 // CHECK-NEXT:     spv.Return
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "Inline" {
      spv.Return
   }
 }
 
 // -----
 
 // CHECK: v1.5
-spv.module Logical GLSL450 requires #spv.vce<v1.5, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.5, [Shader], []> {
 }
 
 // -----
 
 // CHECK: [Shader, Float16]
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader, Float16], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader, Float16], []> {
 }
 
 // -----
 
 // CHECK: [SPV_KHR_float_controls, SPV_KHR_subgroup_vote]
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_float_controls, SPV_KHR_subgroup_vote]> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_float_controls, SPV_KHR_subgroup_vote]> {
 }
 

mlir/test/Target/SPIRV/non-uniform-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @group_non_uniform_ballot
   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_broadcast

mlir/test/Target/SPIRV/ocl-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Physical64 OpenCL requires #spv.vce<v1.0, [Kernel, Addresses], []> {
+spv.mlir.module Physical64 OpenCL requires #spv.vce<v1.0, [Kernel, Addresses], []> {
   spv.func @float_insts(%arg0 : f32) "None" {
     // CHECK: {{%.*}} = spv.OCL.exp {{%.*}} : f32
     %0 = spv.OCL.exp %arg0 : f32
     // CHECK: {{%.*}} = spv.OCL.fabs {{%.*}} : f32
     %1 = spv.OCL.fabs %arg0 : f32
     spv.Return
   }
 

mlir/test/Target/SPIRV/phi.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 // Test branch with one block argument
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
 // Test branch with multiple block arguments
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
 // Test using block arguments within branch
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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: ^bb2(%{{.*}}: i32, %{{.*}}: i32):
   }
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
 // Test block not following domination order
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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:   spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
 // Test multiple predecessors
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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:   spv.mlir.merge
   }
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
 // Test nested loops with block arguments
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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>
   spv.func @fmul_kernel() "None" {
     %3 = spv.Constant 12 : i32
     %4 = spv.Constant 32 : i32
     %5 = spv.Constant 4 : i32
     %6 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %7 = spv.Load "Input" %6 : vector<3xi32>
// CHECK:     spv.mlir.merge
   spv.EntryPoint "GLCompute" @fmul_kernel, @__builtin_var_WorkgroupId__, @__builtin_var_NumWorkgroups__
   spv.ExecutionMode @fmul_kernel "LocalSize", 32, 1, 1
 }
 
 // -----
 
 // Test back-to-back loops with block arguments
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @fmul_kernel() "None" {
     %cst4 = spv.Constant 4 : i32
 
     %val1 = spv.Constant 43 : i32
     %val2 = spv.Constant 44 : i32
 
 // CHECK:        spv.Constant 43
 // CHECK-NEXT:   spv.Branch ^[[BB1:.+]](%{{.+}} : i32)

mlir/test/Target/SPIRV/sampled-image.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: !spv.ptr<!spv.sampled_image<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, UniformConstant>
   spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<!spv.sampled_image<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, UniformConstant>
 
   // CHECK: !spv.ptr<!spv.sampled_image<!spv.image<si32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>>, UniformConstant>
   spv.GlobalVariable @var1 bind(0, 0) : !spv.ptr<!spv.sampled_image<!spv.image<si32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>>, UniformConstant>
 
   // CHECK: !spv.ptr<!spv.sampled_image<!spv.image<i32, Rect, DepthUnknown, Arrayed, MultiSampled, NeedSampler, R8ui>>, UniformConstant>
   spv.GlobalVariable @var2 bind(0, 0) : !spv.ptr<!spv.sampled_image<!spv.image<i32, Rect, DepthUnknown, Arrayed, MultiSampled, NeedSampler, R8ui>>, UniformConstant>
 }

mlir/test/Target/SPIRV/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 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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.mlir.merge
   spv.ExecutionMode @main "LocalSize", 1, 1, 1
 }
 
 // -----
 
 // Selection with only then branch
 // Selection in function entry block
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 // 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

mlir/test/Target/SPIRV/spec-constant.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // 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
 
spv.func @use_composite() -> (i32) "None" {
     %2 = spv.CompositeExtract %0[1 : i32] : !spv.array<2 x i32>
     %3 = spv.IAdd %1, %2 : i32
     spv.ReturnValue %3 : i32
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   spv.SpecConstant @sc_f32_1 = 1.5 : f32
   spv.SpecConstant @sc_f32_2 = 2.5 : f32
   spv.SpecConstant @sc_f32_3 = 3.5 : f32
 
   spv.SpecConstant @sc_i32_1 = 1   : i32
 
   // CHECK: spv.SpecConstantComposite @scc_array (@sc_f32_1, @sc_f32_2, @sc_f32_3) : !spv.array<3 x f32>
   spv.SpecConstantComposite @scc_array (@sc_f32_1, @sc_f32_2, @sc_f32_3) : !spv.array<3 x f32>
 
   // CHECK: spv.SpecConstantComposite @scc_struct (@sc_i32_1, @sc_f32_2, @sc_f32_3) : !spv.struct<(i32, f32, f32)>
   spv.SpecConstantComposite @scc_struct (@sc_i32_1, @sc_f32_2, @sc_f32_3) : !spv.struct<(i32, f32, f32)>
 
   // CHECK: spv.SpecConstantComposite @scc_vector (@sc_f32_1, @sc_f32_2, @sc_f32_3) : vector<3xf32>
   spv.SpecConstantComposite @scc_vector (@sc_f32_1, @sc_f32_2, @sc_f32_3) : vector<3 x f32>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   spv.SpecConstant @sc_f32_1 = 1.5 : f32
   spv.SpecConstant @sc_f32_2 = 2.5 : f32
   spv.SpecConstant @sc_f32_3 = 3.5 : f32
 
   spv.SpecConstant @sc_i32_1 = 1   : i32
 
   // CHECK: spv.SpecConstantComposite @scc_array (@sc_f32_1, @sc_f32_2, @sc_f32_3) : !spv.array<3 x f32>
   spv.SpecConstantComposite @scc_array (@sc_f32_1, @sc_f32_2, @sc_f32_3) : !spv.array<3 x f32>
 
   // CHECK: spv.SpecConstantComposite @scc_struct (@sc_i32_1, @sc_f32_2, @sc_f32_3) : !spv.struct<(i32, f32, f32)>
   spv.SpecConstantComposite @scc_struct (@sc_i32_1, @sc_f32_2, @sc_f32_3) : !spv.struct<(i32, f32, f32)>
 
   // CHECK: spv.SpecConstantComposite @scc_vector (@sc_f32_1, @sc_f32_2, @sc_f32_3) : vector<3xf32>
   spv.SpecConstantComposite @scc_vector (@sc_f32_1, @sc_f32_2, @sc_f32_3) : vector<3 x f32>
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   spv.SpecConstant @sc_i32_1 = 1 : i32
 
   spv.func @use_composite() -> (i32) "None" {
     // CHECK: [[USE1:%.*]] = spv.mlir.referenceof @sc_i32_1 : i32
     // CHECK: [[USE2:%.*]] = spv.Constant 0 : i32
 
     // CHECK: [[RES1:%.*]] = spv.SpecConstantOperation wraps "spv.ISub"([[USE1]], [[USE2]]) : (i32, i32) -> i32

mlir/test/Target/SPIRV/struct.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Input>
   spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Input>
 
   // CHECK: !spv.ptr<!spv.struct<(f32 [0], !spv.struct<(f32 [0], !spv.array<16 x f32, stride=4> [4])> [4])>, Input>
   spv.GlobalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<(f32 [0], !spv.struct<(f32 [0], !spv.array<16 x f32, stride=4> [4])> [4])>, Input>
 
   // CHECK: !spv.ptr<!spv.struct<(f32 [0], i32 [4], f64 [8], i64 [16], f32 [24], i32 [30], f32 [34], i32 [38])>, StorageBuffer>
   spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(f32 [0], i32 [4], f64 [8], i64 [16], f32 [24], i32 [30], f32 [34], i32 [38])>, StorageBuffer>

mlir/test/Target/SPIRV/terminator.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @ret
   spv.func @ret() -> () "None" {
     // CHECK: spv.Return
     spv.Return
   }
 
   // CHECK-LABEL: @ret_val
   spv.func @ret_val() -> (i32) "None" {

mlir/test/Target/SPIRV/undef.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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>, i32
     spv.Return
   }
 }
 
 // -----
 
-spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
+spv.mlir.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: spv.func {{@.*}}
   spv.func @ignore_unused_undef() -> () "None" {
     // CHECK-NEXT: spv.Return
     %0 = spv.undef : f32
     spv.Return
   }
 }

mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp

LogicalResult ConvertToGroupNonUniformBallot::matchAndRewrite(
   Value predicate = op->getOperand(0);
 
   rewriter.replaceOpWithNewOp<spirv::GroupNonUniformBallotOp>(
       op, op->getResult(0).getType(), spirv::Scope::Workgroup, predicate);
   return success();
 }
 
 ConvertToModule::ConvertToModule(MLIRContext *context)
-    : RewritePattern("test.convert_to_module_op", {"spv.module"}, 1, context) {}
+    : RewritePattern("test.convert_to_module_op", {"spv.mlir.module"}, 1,
+                     context) {}
 
 LogicalResult
 ConvertToModule::matchAndRewrite(Operation *op,
                                  PatternRewriter &rewriter) const {
   rewriter.replaceOpWithNewOp<spirv::ModuleOp>(
       op, spirv::AddressingModel::PhysicalStorageBuffer64,
       spirv::MemoryModel::Vulkan);
   return success();

mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp

DeserializationTest() {
     context.getOrLoadDialect<mlir::spirv::SPIRVDialect>();
     // Register a diagnostic handler to capture the diagnostic so that we can
     // check it later.
     context.getDiagEngine().registerHandler([&](Diagnostic &diag) {
       diagnostic.reset(new Diagnostic(std::move(diag)));
     });
   }
 
-  /// Performs deserialization and returns the constructed spv.module op.
+  /// Performs deserialization and returns the constructed spv.mlir.module op.
   spirv::OwningSPIRVModuleRef deserialize() {
     return spirv::deserialize(binary, &context);
   }
 
   /// Checks there is a diagnostic generated with the given `errorMessage`.
   void expectDiagnostic(StringRef errorMessage) {
     ASSERT_NE(nullptr, diagnostic.get());