diff --git a/mlir/docs/Dialects/SPIR-V.md b/mlir/docs/Dialects/SPIR-V.md --- a/mlir/docs/Dialects/SPIR-V.md +++ b/mlir/docs/Dialects/SPIR-V.md @@ -87,7 +87,7 @@ 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 @@ -100,12 +100,12 @@ ## 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. @@ -114,11 +114,11 @@ * 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 @@ -148,7 +148,7 @@ #### 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 @@ -169,7 +169,7 @@ `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`. @@ -200,10 +200,10 @@ * 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 @@ -428,7 +428,7 @@ ``` 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 @@ -960,7 +960,7 @@ 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. @@ -978,7 +978,7 @@ * 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. @@ -1064,7 +1064,7 @@ 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, @@ -1073,7 +1073,7 @@ 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 @@ -1087,7 +1087,7 @@ 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 @@ -1096,7 +1096,7 @@ * [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 @@ -1380,7 +1380,7 @@ `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. diff --git a/mlir/docs/PassManagement.md b/mlir/docs/PassManagement.md --- a/mlir/docs/PassManagement.md +++ b/mlir/docs/PassManagement.md @@ -301,7 +301,7 @@ ``` module { - spv.module "Logical" "GLSL450" { + spv.mlir.module "Logical" "GLSL450" { func @foo() { ... } @@ -313,7 +313,7 @@ ``` `module` - `spv.module` + `spv.mlir.module` `function` ``` diff --git a/mlir/docs/SPIRVToLLVMDialectConversion.md b/mlir/docs/SPIRVToLLVMDialectConversion.md --- a/mlir/docs/SPIRVToLLVMDialectConversion.md +++ b/mlir/docs/SPIRVToLLVMDialectConversion.md @@ -478,7 +478,7 @@ ```mlir // Original SPIR-V module -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @struct : !spv.ptr>, Private> spv.func @func() -> () "None" { %0 = spv.mlir.addressof @struct : !spv.ptr>, Private> @@ -801,13 +801,13 @@ ## 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`. @@ -872,7 +872,7 @@ 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. @@ -896,7 +896,7 @@ 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. diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td --- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td @@ -403,7 +403,7 @@ // ----- -def SPV_ModuleOp : SPV_Op<"module", +def SPV_ModuleOp : SPV_Op<"mlir.module", [IsolatedFromAbove, SingleBlockImplicitTerminator<"ModuleEndOp">, SymbolTable, Symbol]> { @@ -430,7 +430,7 @@ ``` 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 @@ -439,9 +439,9 @@ #### Example: ```mlir - spv.module Logical GLSL450 {} + spv.mlir.module Logical GLSL450 {} - spv.module Logical Vulkan + spv.mlir.module Logical Vulkan requires #spv.vce attributes { some_additional_attr = ... } { spv.func @do_nothing() -> () { @@ -497,7 +497,7 @@ 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. diff --git a/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h b/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h --- a/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.h @@ -32,13 +32,13 @@ 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> createUpdateVersionCapabilityExtensionPass(); diff --git a/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td b/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td --- a/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td +++ b/mlir/include/mlir/Dialect/SPIRV/Transforms/Passes.td @@ -30,7 +30,7 @@ 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()"; } diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp --- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp +++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp @@ -62,7 +62,8 @@ 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 { public: using OpConversionPattern::OpConversionPattern; @@ -75,7 +76,7 @@ SmallVector 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 { public: using OpConversionPattern::OpConversionPattern; @@ -299,8 +300,8 @@ 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); diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp --- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp +++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // // 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. // //===----------------------------------------------------------------------===// diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp --- a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp +++ b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp @@ -134,7 +134,8 @@ SmallVector binary; for (auto spirvModule : module.getOps()) { 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))) diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp --- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp +++ b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp @@ -48,7 +48,7 @@ target.addIllegalDialect(); target.addLegalDialect(); - // Set `ModuleOp` and `ModuleTerminatorOp` as legal for `spv.module` + // Set `ModuleOp` and `ModuleTerminatorOp` as legal for `spv.mlir.module` // conversion. target.addLegalOp(); target.addLegalOp(); diff --git a/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp b/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp --- a/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp +++ b/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp @@ -234,8 +234,8 @@ /// 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 { public: using OpConversionPattern::OpConversionPattern; diff --git a/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp b/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp --- a/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp +++ b/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp @@ -2473,7 +2473,7 @@ } //===----------------------------------------------------------------------===// -// spv.module +// spv.mlir.module //===----------------------------------------------------------------------===// void spirv::ModuleOp::build(OpBuilder &builder, OperationState &state, @@ -2570,7 +2570,7 @@ 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 @@ -2579,7 +2579,7 @@ auto funcOp = table.lookup(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) { @@ -2609,14 +2609,15 @@ entryPoints[key] = entryPointOp; } else if (auto funcOp = dyn_cast(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"); } } } diff --git a/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp b/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp --- a/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp +++ b/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp @@ -147,7 +147,7 @@ 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: diff --git a/mlir/lib/Target/SPIRV/Serialization/SerializeOps.cpp b/mlir/lib/Target/SPIRV/Serialization/SerializeOps.cpp --- a/mlir/lib/Target/SPIRV/Serialization/SerializeOps.cpp +++ b/mlir/lib/Target/SPIRV/Serialization/SerializeOps.cpp @@ -552,9 +552,10 @@ for (auto var : interface.getValue()) { auto id = getVariableID(var.cast().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 " + "referenced variables should already be defined"); } operands.push_back(id); } diff --git a/mlir/lib/Target/SPIRV/TranslateRegistration.cpp b/mlir/lib/Target/SPIRV/TranslateRegistration.cpp --- a/mlir/lib/Target/SPIRV/TranslateRegistration.cpp +++ b/mlir/lib/Target/SPIRV/TranslateRegistration.cpp @@ -89,10 +89,10 @@ 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))) @@ -128,10 +128,10 @@ auto spirvModules = srcModule.getOps(); 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(); diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir --- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir @@ -8,7 +8,7 @@ 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 @@ -32,7 +32,7 @@ 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 @@ -56,7 +56,7 @@ 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 @@ -80,7 +80,7 @@ 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>}} { @@ -105,7 +105,7 @@ 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>}} { @@ -127,7 +127,7 @@ 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>}} { @@ -149,7 +149,7 @@ 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 @@ -173,7 +173,7 @@ 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 @@ -190,7 +190,7 @@ // ----- 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 @@ -206,7 +206,7 @@ // ----- 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 @@ -222,7 +222,7 @@ // ----- 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 diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir --- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir @@ -22,7 +22,7 @@ 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, Input> // CHECK-DAG: spv.GlobalVariable @[[$LOCALINVOCATIONIDVAR:.*]] built_in("LocalInvocationId") : !spv.ptr, Input> diff --git a/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir b/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir --- a/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir @@ -5,7 +5,7 @@ spv.target_env = #spv.target_env<#spv.vce, {}> } { 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 diff --git a/mlir/test/Conversion/GPUToSPIRV/simple.mlir b/mlir/test/Conversion/GPUToSPIRV/simple.mlir --- a/mlir/test/Conversion/GPUToSPIRV/simple.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/simple.mlir @@ -2,7 +2,7 @@ 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 [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>} @@ -29,7 +29,7 @@ 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> diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir --- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir +++ b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir @@ -5,7 +5,7 @@ // 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 { + spv.mlir.module Logical GLSL450 requires #spv.vce { spv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr [0])>, StorageBuffer> spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} { %0 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr [0])>, StorageBuffer> diff --git a/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir @@ -4,7 +4,7 @@ // spv.Branch //===----------------------------------------------------------------------===// -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @branch_without_arguments() -> () "None" { // CHECK: llvm.br ^bb1 spv.Branch ^label @@ -30,7 +30,7 @@ // 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 @@ -85,7 +85,7 @@ // 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:.*]] @@ -124,7 +124,7 @@ // spv.selection //===----------------------------------------------------------------------===// -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @selection_empty() -> () "None" { // CHECK: llvm.return spv.selection { diff --git a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir @@ -5,7 +5,7 @@ // 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 [0])>, StorageBuffer> // CHECK: spv.func @__spv__foo_bar @@ -21,7 +21,7 @@ // CHECK-NEXT: llvm.mlir.constant(false) : i1 // CHECK-NEXT: "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr, !llvm.ptr)>>, i64, i1) -> () - spv.module @__spv__foo Logical GLSL450 requires #spv.vce { + spv.mlir.module @__spv__foo Logical GLSL450 requires #spv.vce { spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr [0])>, StorageBuffer> spv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} { %0 = spv.mlir.addressof @bar_arg_0 : !spv.ptr [0])>, StorageBuffer> diff --git a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir @@ -29,12 +29,12 @@ // 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 } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // CHECK: llvm.mlir.global private @struct() : !llvm.struct)> // CHECK-LABEL: @func // CHECK: llvm.mlir.addressof @struct : !llvm.ptr)>> @@ -45,7 +45,7 @@ } } -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 @@ -56,7 +56,7 @@ } } -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 diff --git a/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir @@ -73,7 +73,7 @@ // CHECK-NEXT: llvm.return // CHECK-NEXT: } // CHECK-NEXT: } -spv.module Logical OpenCL { +spv.mlir.module Logical OpenCL { spv.func @empty() "None" { spv.Return } @@ -98,7 +98,7 @@ // CHECK-NEXT: llvm.return // CHECK-NEXT: } // CHECK-NEXT: } -spv.module Logical OpenCL { +spv.mlir.module Logical OpenCL { spv.func @bar() "None" { spv.Return } diff --git a/mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir @@ -1,26 +1,26 @@ // 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 {} +spv.mlir.module Logical GLSL450 requires #spv.vce {} // 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 diff --git a/mlir/test/Dialect/SPIRV/IR/availability.mlir b/mlir/test/Dialect/SPIRV/IR/availability.mlir --- a/mlir/test/Dialect/SPIRV/IR/availability.mlir +++ b/mlir/test/Dialect/SPIRV/IR/availability.mlir @@ -32,20 +32,20 @@ // 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 } diff --git a/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir b/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir @@ -155,7 +155,7 @@ // 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> @@ -200,7 +200,7 @@ // ----- -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) @@ -210,7 +210,7 @@ // ----- -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) @@ -220,7 +220,7 @@ // ----- -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) -> () @@ -230,7 +230,7 @@ // ----- -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}} @@ -241,7 +241,7 @@ // ----- -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'}} @@ -252,7 +252,7 @@ // ----- -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 @@ -528,7 +528,7 @@ // ----- // 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 @@ -537,7 +537,7 @@ // ----- -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 @@ -615,7 +615,7 @@ // ----- -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}} @@ -625,7 +625,7 @@ // ----- -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')}} @@ -635,7 +635,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @in_nested_region(%cond: i1) -> () "None" { spv.selection { spv.BranchConditional %cond, ^then, ^merge diff --git a/mlir/test/Dialect/SPIRV/IR/matrix-ops.mlir b/mlir/test/Dialect/SPIRV/IR/matrix-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/matrix-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/matrix-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -verify-diagnostics %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // 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>> diff --git a/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir b/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir @@ -339,7 +339,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @var0 : !spv.ptr // CHECK_LABEL: @simple_load spv.func @simple_load() -> () "None" { @@ -462,7 +462,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @var0 : !spv.ptr spv.func @simple_store(%arg0 : f32) -> () "None" { %0 = spv.mlir.addressof @var0 : !spv.ptr @@ -495,7 +495,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @global : !spv.ptr spv.func @variable_init_global_variable() -> () "None" { %0 = spv.mlir.addressof @global : !spv.ptr @@ -507,7 +507,7 @@ // ----- -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" { diff --git a/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir b/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir @@ -4,7 +4,7 @@ // spv.mlir.addressof //===----------------------------------------------------------------------===// -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @var1 : !spv.ptr)>, Input> spv.func @access_chain() -> () "None" { %0 = spv.Constant 1: i32 @@ -28,7 +28,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @var1 : !spv.ptr)>, Input> spv.func @foo() -> () "None" { // expected-error @+1 {{expected spv.GlobalVariable symbol}} @@ -38,7 +38,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @var1 : !spv.ptr)>, Input> spv.func @foo() -> () "None" { // expected-error @+1 {{result type mismatch with the referenced global variable's type}} @@ -135,7 +135,7 @@ // spv.EntryPoint //===----------------------------------------------------------------------===// -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -143,7 +143,7 @@ spv.EntryPoint "GLCompute" @do_nothing } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @var2 : !spv.ptr spv.GlobalVariable @var3 : !spv.ptr spv.func @do_something(%arg0 : !spv.ptr, %arg1 : !spv.ptr) -> () "None" { @@ -157,7 +157,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -167,11 +167,11 @@ // ----- -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 } @@ -182,7 +182,7 @@ // ----- -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 @@ -191,7 +191,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -202,7 +202,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -217,7 +217,7 @@ // spv.ExecutionMode //===----------------------------------------------------------------------===// -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -226,7 +226,7 @@ spv.ExecutionMode @do_nothing "ContractionOff" } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -237,7 +237,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @do_nothing() -> () "None" { spv.Return } @@ -285,7 +285,7 @@ // ----- // 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" { @@ -301,13 +301,13 @@ // spv.GlobalVariable //===----------------------------------------------------------------------===// -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // CHECK: spv.GlobalVariable @var0 : !spv.ptr spv.GlobalVariable @var0 : !spv.ptr } // 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 // spv.GlobalVariable @var1 init(%0) : !spv.ptr @@ -315,7 +315,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.SpecConstant @sc = 4.0 : f32 // CHECK: spv.GlobalVariable @var initializer(@sc) : !spv.ptr spv.GlobalVariable @var initializer(@sc) : !spv.ptr @@ -330,13 +330,13 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // CHECK: spv.GlobalVariable @var0 bind(1, 2) : !spv.ptr spv.GlobalVariable @var0 bind(1, 2) : !spv.ptr } // 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 // spv.GlobalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr @@ -344,7 +344,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // CHECK: spv.GlobalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr, Input> spv.GlobalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr, Input> // CHECK: spv.GlobalVariable @var2 built_in("GlobalInvocationID") : !spv.ptr, Input> @@ -361,35 +361,35 @@ // ----- -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 } // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // expected-error @+1 {{storage class cannot be 'Generic'}} spv.GlobalVariable @var0 : !spv.ptr } // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // expected-error @+1 {{storage class cannot be 'Function'}} spv.GlobalVariable @var0 : !spv.ptr } // ----- -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 @@ -400,40 +400,40 @@ // ----- //===----------------------------------------------------------------------===// -// 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 -spv.module Logical GLSL450 requires #spv.vce { } +// CHECK: spv.mlir.module Logical GLSL450 requires #spv.vce +spv.mlir.module Logical GLSL450 requires #spv.vce { } // 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 attributes {foo = "bar"} -spv.module Logical GLSL450 +// CHECK: spv.mlir.module Logical GLSL450 requires #spv.vce attributes {foo = "bar"} +spv.mlir.module Logical GLSL450 requires #spv.vce 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 } @@ -442,32 +442,32 @@ // ----- // 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: @@ -479,24 +479,24 @@ // 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"() : () -> () } } @@ -504,8 +504,8 @@ // ----- // 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" } @@ -526,7 +526,7 @@ // 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 @@ -591,7 +591,7 @@ // ----- -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 @@ -601,7 +601,7 @@ // ----- -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}} @@ -612,7 +612,7 @@ // ----- -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" { @@ -628,7 +628,7 @@ // 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 @@ -639,21 +639,21 @@ // ----- -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> } @@ -672,7 +672,7 @@ // 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 } @@ -683,7 +683,7 @@ // ----- -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 @@ -693,7 +693,7 @@ // ----- -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 @@ -704,7 +704,7 @@ // ----- -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 @@ -718,7 +718,7 @@ // ----- -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 @@ -728,7 +728,7 @@ // ----- -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 @@ -738,7 +738,7 @@ // ----- -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 @@ -752,7 +752,7 @@ // ----- -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 @@ -762,7 +762,7 @@ // ----- -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 @@ -773,7 +773,7 @@ // ----- -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 @@ -787,7 +787,7 @@ // ----- -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> @@ -799,7 +799,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo() -> i32 "None" { // CHECK: [[LHS:%.*]] = spv.Constant %0 = spv.Constant 1: i32 @@ -815,7 +815,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.SpecConstant @sc = 42 : i32 spv.func @foo() -> i32 "None" { @@ -829,7 +829,7 @@ // ----- -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'}} @@ -839,7 +839,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo() -> () "None" { %0 = spv.Variable : !spv.ptr @@ -851,7 +851,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo() -> () "None" { %0 = spv.Variable : !spv.ptr %1 = spv.Load "Function" %0 : i32 diff --git a/mlir/test/Dialect/SPIRV/IR/target-env.mlir b/mlir/test/Dialect/SPIRV/IR/target-env.mlir --- a/mlir/test/Dialect/SPIRV/IR/target-env.mlir +++ b/mlir/test/Dialect/SPIRV/IR/target-env.mlir @@ -148,7 +148,7 @@ func @module_suitable_extension1() attributes { spv.target_env = #spv.target_env<#spv.vce, {}> } { - // CHECK: spv.module PhysicalStorageBuffer64 Vulkan + // CHECK: spv.mlir.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () ->() return } @@ -157,7 +157,7 @@ func @module_suitable_extension2() attributes { spv.target_env = #spv.target_env<#spv.vce, {}> } { - // CHECK: spv.module PhysicalStorageBuffer64 Vulkan + // CHECK: spv.mlir.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () -> () return } @@ -185,7 +185,7 @@ // Version 1.5 implies SPV_KHR_vulkan_memory_model and SPV_KHR_physical_storage_buffer. spv.target_env = #spv.target_env<#spv.vce, {}> } { - // CHECK: spv.module PhysicalStorageBuffer64 Vulkan + // CHECK: spv.mlir.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () -> () return } diff --git a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir --- a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir +++ b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir @@ -1,7 +1,7 @@ // 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 @@ -13,11 +13,11 @@ // 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 @@ -30,21 +30,21 @@ // ----- 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 { } } diff --git a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict-resolution.mlir b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict-resolution.mlir --- a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict-resolution.mlir +++ b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict-resolution.mlir @@ -3,7 +3,7 @@ // 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: } @@ -15,13 +15,13 @@ // 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 } @@ -33,7 +33,7 @@ // 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: } @@ -51,20 +51,20 @@ // 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 @@ -77,7 +77,7 @@ // 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: } @@ -94,13 +94,13 @@ // 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 } @@ -118,7 +118,7 @@ // 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: } @@ -135,13 +135,13 @@ // 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 @@ -159,7 +159,7 @@ // 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: } @@ -174,13 +174,13 @@ // 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 } @@ -193,7 +193,7 @@ // ----- // 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: } @@ -211,7 +211,7 @@ // CHECK-NEXT: } module { -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo(%arg0 : i32) -> i32 "None" { spv.ReturnValue %arg0 : i32 } @@ -220,7 +220,7 @@ spv.ExecutionMode @foo "ContractionOff" } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo(%arg0 : f32) -> f32 "None" { spv.ReturnValue %arg0 : f32 } @@ -235,7 +235,7 @@ // 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: } @@ -244,13 +244,13 @@ // 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 } } @@ -261,7 +261,7 @@ // 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: } @@ -275,13 +275,13 @@ // 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 spv.func @bar() -> f32 "None" { @@ -298,7 +298,7 @@ // 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 @@ -312,7 +312,7 @@ // CHECK-NEXT: } module { -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @foo bind(1, 0) : !spv.ptr spv.func @bar() -> f32 "None" { @@ -322,7 +322,7 @@ } } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo(%arg0 : i32) -> i32 "None" { spv.ReturnValue %arg0 : i32 } @@ -334,7 +334,7 @@ // 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: } @@ -343,13 +343,13 @@ // 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 } } @@ -360,7 +360,7 @@ // 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: } @@ -373,13 +373,13 @@ // 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" { @@ -395,7 +395,7 @@ // 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 @@ -408,7 +408,7 @@ // CHECK-NEXT: } module { -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.SpecConstant @foo = -5 : i32 spv.func @bar() -> i32 "None" { @@ -417,7 +417,7 @@ } } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo(%arg0 : i32) -> i32 "None" { spv.ReturnValue %arg0 : i32 } @@ -429,7 +429,7 @@ // 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: } @@ -439,13 +439,13 @@ // 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> } @@ -457,7 +457,7 @@ // 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: } @@ -472,13 +472,13 @@ // 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> @@ -496,7 +496,7 @@ // 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 @@ -511,7 +511,7 @@ // 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> @@ -522,7 +522,7 @@ } } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo(%arg0 : i32) -> i32 "None" { spv.ReturnValue %arg0 : i32 } @@ -535,7 +535,7 @@ // 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 @@ -554,7 +554,7 @@ // 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> @@ -565,7 +565,7 @@ } } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @foo(%arg0 : i32) -> i32 "None" { spv.ReturnValue %arg0 : i32 } @@ -581,18 +581,18 @@ // 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 } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @foo bind(2, 0) : !spv.ptr } } @@ -600,18 +600,18 @@ // ----- // 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, Input> } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @foo built_in("LocalInvocationId") : !spv.ptr, Input> } } @@ -621,18 +621,18 @@ // 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 } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.SpecConstant @foo = -5 : i32 } } @@ -642,18 +642,18 @@ // 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 } } @@ -663,7 +663,7 @@ // 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 @@ -671,11 +671,11 @@ // CHECK-NEXT: } module { -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @foo bind(1, 0) : !spv.ptr } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.SpecConstant @bar = -5 : i32 spv.SpecConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32> } @@ -686,7 +686,7 @@ // 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) @@ -694,12 +694,12 @@ // 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 } } diff --git a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir --- a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir +++ b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir @@ -3,7 +3,7 @@ // 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 @@ -22,7 +22,7 @@ // CHECK-NEXT: } module { -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @foo bind(1, 0) : !spv.ptr spv.func @use_foo() -> f32 "None" { @@ -32,7 +32,7 @@ } } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @bar bind(1, 0) : !spv.ptr spv.func @use_bar() -> f32 "None" { @@ -49,7 +49,7 @@ // 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) @@ -63,11 +63,11 @@ // CHECK-NEXT: } module { -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @foo bind(1, 0) : !spv.ptr } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @bar bind(1, 0) : !spv.ptr spv.func @use_bar() -> f32 "None" { @@ -83,7 +83,7 @@ // 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 @@ -94,11 +94,11 @@ // CHECK-NEXT: } module { -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @foo built_in("GlobalInvocationId") : !spv.ptr, Input> } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @bar built_in("GlobalInvocationId") : !spv.ptr, Input> spv.func @use_bar() -> vector<3xi32> "None" { @@ -112,7 +112,7 @@ // ----- // 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() @@ -129,7 +129,7 @@ // 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" { @@ -138,7 +138,7 @@ } } -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.SpecConstant @bar spec_id(5) = 1. : f32 spv.func @use_bar() -> (f32) "None" { @@ -152,7 +152,7 @@ // ----- // 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) @@ -191,7 +191,7 @@ // 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" { diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir @@ -3,8 +3,8 @@ module attributes { spv.target_env = #spv.target_env<#spv.vce, {}> } { - 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)>, CrossWorkgroup> // CHECK: spv.EntryPoint "Kernel" [[FN]] // CHECK: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir @@ -5,8 +5,8 @@ #spv.vce, {}> } { -// 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, StorageBuffer> // CHECK-DAG: spv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr [0])>, StorageBuffer> // CHECK: spv.func [[FN:@.*]]() @@ -26,6 +26,6 @@ } // CHECK: spv.EntryPoint "GLCompute" [[FN]] // CHECK: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 -} // end spv.module +} // end spv.mlir.module } // end module diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir @@ -5,8 +5,8 @@ #spv.vce, {}> } { -// 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, Input> // CHECK-DAG: spv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") @@ -116,6 +116,6 @@ } // 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 diff --git a/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir b/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir @@ -1,6 +1,6 @@ -// 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 } @@ -15,7 +15,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @callee() -> i32 "None" { %0 = spv.Constant 42 : i32 spv.ReturnValue %0 : i32 @@ -32,7 +32,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @data bind(0, 0) : !spv.ptr [0])>, StorageBuffer> spv.func @callee() "None" { %0 = spv.mlir.addressof @data : !spv.ptr [0])>, StorageBuffer> @@ -67,7 +67,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @callee(%cond : i1) -> () "None" { spv.selection { spv.BranchConditional %cond, ^then, ^merge @@ -90,7 +90,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @callee(%cond : i1) -> () "None" { spv.selection { spv.BranchConditional %cond, ^then, ^merge @@ -119,7 +119,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @callee(%cond : i1) -> () "None" { spv.loop { spv.Branch ^header @@ -146,7 +146,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.func @callee(%cond : i1) -> () "None" { spv.loop { spv.Branch ^header @@ -183,7 +183,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { spv.GlobalVariable @arg_0 bind(0, 0) : !spv.ptr, StorageBuffer> spv.GlobalVariable @arg_1 bind(0, 1) : !spv.ptr, StorageBuffer> diff --git a/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir b/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir @@ -1,6 +1,6 @@ // 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 [4], f32 [12])>, Uniform> spv.GlobalVariable @var0 bind(0,1) : !spv.ptr, f32)>, Uniform> @@ -31,7 +31,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // CHECK: spv.GlobalVariable @var0 : !spv.ptr [0], i1 [16])> [0], i1 [24])> [0], i1 [32])> [0], i1 [40])>, Uniform> spv.GlobalVariable @var0 : !spv.ptr, i1)>, i1)>, i1)>, i1)>, Uniform> @@ -59,7 +59,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // CHECK: spv.GlobalVariable @var0 : !spv.ptr [0], f32 [8])>, StorageBuffer> spv.GlobalVariable @var0 : !spv.ptr, f32)>, StorageBuffer> @@ -72,7 +72,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // CHECK: spv.GlobalVariable @emptyStructAsMember : !spv.ptr [0])>, StorageBuffer> spv.GlobalVariable @emptyStructAsMember : !spv.ptr)>, StorageBuffer> @@ -91,7 +91,7 @@ // ----- -spv.module Logical GLSL450 { +spv.mlir.module Logical GLSL450 { // CHECK: spv.GlobalVariable @var0 : !spv.ptr, PushConstant> spv.GlobalVariable @var0 : !spv.ptr, PushConstant> // CHECK: spv.GlobalVariable @var1 : !spv.ptr, PhysicalStorageBuffer> diff --git a/mlir/test/Dialect/SPIRV/Transforms/rewrite-inserts.mlir b/mlir/test/Dialect/SPIRV/Transforms/rewrite-inserts.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/rewrite-inserts.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/rewrite-inserts.mlir @@ -1,6 +1,6 @@ // 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> diff --git a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir @@ -8,7 +8,7 @@ // spv.IAdd is available from v1.0. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -22,7 +22,7 @@ // spv.GroupNonUniformBallot is available since v1.3. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -39,7 +39,7 @@ // Test minimal capabilities. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -53,7 +53,7 @@ // AtomicStorage implies Shader. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -74,7 +74,7 @@ // * GroupNonUniformBallot // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -85,7 +85,7 @@ } // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -99,7 +99,7 @@ // Using 8-bit integers in non-interface storage class requires Int8. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -111,7 +111,7 @@ // Using 16-bit floats in non-interface storage class requires Float16. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -123,7 +123,7 @@ // Using 16-element vectors requires Vector16. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -141,7 +141,7 @@ // spv.SubgroupBallotKHR requires the SPV_KHR_shader_ballot extension. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> @@ -157,7 +157,7 @@ // implicitly by v1.5. // CHECK: requires #spv.vce -spv.module Logical Vulkan attributes { +spv.mlir.module Logical Vulkan attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -172,7 +172,7 @@ // Using 8-bit integers in interface storage class requires additional // extensions and capabilities. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> } { @@ -187,7 +187,7 @@ // * Buffer requires ImageBuffer or SampledBuffer. // * Rg32f requires StorageImageExtendedFormats. // CHECK: requires #spv.vce -spv.module Logical GLSL450 attributes { +spv.mlir.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> diff --git a/mlir/test/Target/SPIRV/arithmetic-ops.mlir b/mlir/test/Target/SPIRV/arithmetic-ops.mlir --- a/mlir/test/Target/SPIRV/arithmetic-ops.mlir +++ b/mlir/test/Target/SPIRV/arithmetic-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @fmul(%arg0 : f32, %arg1 : f32) "None" { // CHECK: {{%.*}}= spv.FMul {{%.*}}, {{%.*}} : f32 %0 = spv.FMul %arg0, %arg1 : f32 diff --git a/mlir/test/Target/SPIRV/array.mlir b/mlir/test/Target/SPIRV/array.mlir --- a/mlir/test/Target/SPIRV/array.mlir +++ b/mlir/test/Target/SPIRV/array.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @array_stride(%arg0 : !spv.ptr, stride=128>, StorageBuffer>, %arg1 : i32, %arg2 : i32) "None" { // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr, stride=128>, StorageBuffer>, i32, i32 %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr, stride=128>, StorageBuffer>, i32, i32 @@ -10,7 +10,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr, StorageBuffer> spv.GlobalVariable @var0 : !spv.ptr, StorageBuffer> // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr>, Input> diff --git a/mlir/test/Target/SPIRV/atomic-ops.mlir b/mlir/test/Target/SPIRV/atomic-ops.mlir --- a/mlir/test/Target/SPIRV/atomic-ops.mlir +++ b/mlir/test/Target/SPIRV/atomic-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @atomic_compare_exchange_weak spv.func @atomic_compare_exchange_weak(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 "None" { // CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %{{.*}}, %{{.*}}, %{{.*}} : !spv.ptr diff --git a/mlir/test/Target/SPIRV/barrier-ops.mlir b/mlir/test/Target/SPIRV/barrier-ops.mlir --- a/mlir/test/Target/SPIRV/barrier-ops.mlir +++ b/mlir/test/Target/SPIRV/barrier-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @memory_barrier_0() -> () "None" { // CHECK: spv.MemoryBarrier Device, "Release|UniformMemory" spv.MemoryBarrier Device, "Release|UniformMemory" diff --git a/mlir/test/Target/SPIRV/bit-ops.mlir b/mlir/test/Target/SPIRV/bit-ops.mlir --- a/mlir/test/Target/SPIRV/bit-ops.mlir +++ b/mlir/test/Target/SPIRV/bit-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @bitcount(%arg: i32) -> i32 "None" { // CHECK: spv.BitCount {{%.*}} : i32 %0 = spv.BitCount %arg : i32 diff --git a/mlir/test/Target/SPIRV/cast-ops.mlir b/mlir/test/Target/SPIRV/cast-ops.mlir --- a/mlir/test/Target/SPIRV/cast-ops.mlir +++ b/mlir/test/Target/SPIRV/cast-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @bit_cast(%arg0 : f32) "None" { // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : f32 to i32 %0 = spv.Bitcast %arg0 : f32 to i32 @@ -14,7 +14,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @convert_f_to_s(%arg0 : f32) -> i32 "None" { // CHECK: {{%.*}} = spv.ConvertFToS {{%.*}} : f32 to i32 %0 = spv.ConvertFToS %arg0 : f32 to i32 diff --git a/mlir/test/Target/SPIRV/composite-op.mlir b/mlir/test/Target/SPIRV/composite-op.mlir --- a/mlir/test/Target/SPIRV/composite-op.mlir +++ b/mlir/test/Target/SPIRV/composite-op.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { 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)>)> diff --git a/mlir/test/Target/SPIRV/constant.mlir b/mlir/test/Target/SPIRV/constant.mlir --- a/mlir/test/Target/SPIRV/constant.mlir +++ b/mlir/test/Target/SPIRV/constant.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @bool_const spv.func @bool_const() -> () "None" { // CHECK: spv.Constant true diff --git a/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir b/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir --- a/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir +++ b/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @cooperative_matrix_load spv.func @cooperative_matrix_load(%ptr : !spv.ptr, %stride : i32, %b : i1) "None" { // CHECK: {{%.*}} = spv.CooperativeMatrixLoadNV {{%.*}}, {{%.*}}, {{%.*}} : !spv.ptr as !spv.coopmatrix<16x8xi32, Workgroup> diff --git a/mlir/test/Target/SPIRV/debug.mlir b/mlir/test/Target/SPIRV/debug.mlir --- a/mlir/test/Target/SPIRV/debug.mlir +++ b/mlir/test/Target/SPIRV/debug.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip-debug -mlir-print-debuginfo -mlir-print-local-scope %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: loc({{".*debug.mlir"}}:5:3) spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr spv.func @arithmetic(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" { diff --git a/mlir/test/Target/SPIRV/decorations.mlir b/mlir/test/Target/SPIRV/decorations.mlir --- a/mlir/test/Target/SPIRV/decorations.mlir +++ b/mlir/test/Target/SPIRV/decorations.mlir @@ -1,27 +1,27 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: location = 0 : i32 spv.GlobalVariable @var {location = 0 : i32} : !spv.ptr, Input> } // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: no_perspective spv.GlobalVariable @var {no_perspective} : !spv.ptr, Input> } // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: flat spv.GlobalVariable @var {flat} : !spv.ptr } // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: aliased // CHECK: aliased spv.GlobalVariable @var1 bind(0, 0) {aliased} : !spv.ptr[0])>, StorageBuffer> @@ -30,21 +30,21 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: non_readable spv.GlobalVariable @var bind(0, 0) {non_readable} : !spv.ptr[0])>, StorageBuffer> } // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: non_writable spv.GlobalVariable @var bind(0, 0) {non_writable} : !spv.ptr[0])>, StorageBuffer> } // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: restrict spv.GlobalVariable @var bind(0, 0) {restrict} : !spv.ptr[0])>, StorageBuffer> } diff --git a/mlir/test/Target/SPIRV/entry-point.mlir b/mlir/test/Target/SPIRV/entry-point.mlir --- a/mlir/test/Target/SPIRV/entry-point.mlir +++ b/mlir/test/Target/SPIRV/entry-point.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @noop() -> () "None" { spv.Return } @@ -12,7 +12,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: spv.GlobalVariable @var2 : !spv.ptr // CHECK-NEXT: spv.GlobalVariable @var3 : !spv.ptr // CHECK-NEXT: spv.func @noop({{%.*}}: !spv.ptr, {{%.*}}: !spv.ptr) "None" diff --git a/mlir/test/Target/SPIRV/execution-mode.mlir b/mlir/test/Target/SPIRV/execution-mode.mlir --- a/mlir/test/Target/SPIRV/execution-mode.mlir +++ b/mlir/test/Target/SPIRV/execution-mode.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { spv.Return } diff --git a/mlir/test/Target/SPIRV/function-call.mlir b/mlir/test/Target/SPIRV/function-call.mlir --- a/mlir/test/Target/SPIRV/function-call.mlir +++ b/mlir/test/Target/SPIRV/function-call.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.GlobalVariable @var1 : !spv.ptr, Input> spv.func @fmain() -> i32 "None" { %0 = spv.Constant 16 : i32 diff --git a/mlir/test/Target/SPIRV/global-variable.mlir b/mlir/test/Target/SPIRV/global-variable.mlir --- a/mlir/test/Target/SPIRV/global-variable.mlir +++ b/mlir/test/Target/SPIRV/global-variable.mlir @@ -5,7 +5,7 @@ // CHECK-NEXT: spv.GlobalVariable @var2 built_in("GlobalInvocationId") : !spv.ptr, Input> // CHECK-NEXT: spv.GlobalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr, Input> -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.GlobalVariable @var0 bind(1, 0) : !spv.ptr spv.GlobalVariable @var1 bind(0, 1) : !spv.ptr spv.GlobalVariable @var2 {built_in = "GlobalInvocationId"} : !spv.ptr, Input> @@ -14,7 +14,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: spv.GlobalVariable @var1 : !spv.ptr // CHECK-NEXT: spv.GlobalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr spv.GlobalVariable @var1 : !spv.ptr @@ -23,7 +23,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.GlobalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr, Input> spv.func @foo() "None" { // CHECK: %[[ADDR:.*]] = spv.mlir.addressof @globalInvocationID : !spv.ptr, Input> diff --git a/mlir/test/Target/SPIRV/glsl-ops.mlir b/mlir/test/Target/SPIRV/glsl-ops.mlir --- a/mlir/test/Target/SPIRV/glsl-ops.mlir +++ b/mlir/test/Target/SPIRV/glsl-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @fmul(%arg0 : f32, %arg1 : f32, %arg2 : i32) "None" { // CHECK: {{%.*}} = spv.GLSL.Exp {{%.*}} : f32 %0 = spv.GLSL.Exp %arg0 : f32 diff --git a/mlir/test/Target/SPIRV/group-ops.mlir b/mlir/test/Target/SPIRV/group-ops.mlir --- a/mlir/test/Target/SPIRV/group-ops.mlir +++ b/mlir/test/Target/SPIRV/group-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @subgroup_ballot spv.func @subgroup_ballot(%predicate: i1) -> vector<4xi32> "None" { // CHECK: %{{.*}} = spv.SubgroupBallotKHR %{{.*}}: vector<4xi32> diff --git a/mlir/test/Target/SPIRV/image.mlir b/mlir/test/Target/SPIRV/image.mlir --- a/mlir/test/Target/SPIRV/image.mlir +++ b/mlir/test/Target/SPIRV/image.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: !spv.ptr, UniformConstant> spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr, UniformConstant> diff --git a/mlir/test/Target/SPIRV/logical-ops.mlir b/mlir/test/Target/SPIRV/logical-ops.mlir --- a/mlir/test/Target/SPIRV/logical-ops.mlir +++ b/mlir/test/Target/SPIRV/logical-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @iequal_scalar(%arg0: i32, %arg1: i32) "None" { // CHECK: {{.*}} = spv.IEqual {{.*}}, {{.*}} : i32 %0 = spv.IEqual %arg0, %arg1 : i32 @@ -90,7 +90,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.SpecConstant @condition_scalar = true spv.func @select() -> () "None" { %0 = spv.Constant 4.0 : f32 diff --git a/mlir/test/Target/SPIRV/loop.mlir b/mlir/test/Target/SPIRV/loop.mlir --- a/mlir/test/Target/SPIRV/loop.mlir +++ b/mlir/test/Target/SPIRV/loop.mlir @@ -2,7 +2,7 @@ // Single loop -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // for (int i = 0; i < count; ++i) {} spv.func @loop(%count : i32) -> () "None" { %zero = spv.Constant 0: i32 @@ -59,7 +59,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.GlobalVariable @GV1 bind(0, 0) : !spv.ptr [0])>, StorageBuffer> spv.GlobalVariable @GV2 bind(0, 1) : !spv.ptr [0])>, StorageBuffer> spv.func @loop_kernel() "None" { @@ -107,7 +107,7 @@ // Nested loop -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // for (int i = 0; i < count; ++i) { // for (int j = 0; j < count; ++j) { } // } diff --git a/mlir/test/Target/SPIRV/matrix.mlir b/mlir/test/Target/SPIRV/matrix.mlir --- a/mlir/test/Target/SPIRV/matrix.mlir +++ b/mlir/test/Target/SPIRV/matrix.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @matrix_access_chain spv.func @matrix_access_chain(%arg0 : !spv.ptr>, Function>, %arg1 : i32) -> !spv.ptr, Function> "None" { // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr>, Function> @@ -47,7 +47,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr>, StorageBuffer> spv.GlobalVariable @var0 : !spv.ptr>, StorageBuffer> diff --git a/mlir/test/Target/SPIRV/memory-ops.mlir b/mlir/test/Target/SPIRV/memory-ops.mlir --- a/mlir/test/Target/SPIRV/memory-ops.mlir +++ b/mlir/test/Target/SPIRV/memory-ops.mlir @@ -4,7 +4,7 @@ // CHECK-NEXT: [[VALUE:%.*]] = spv.Load "Input" [[ARG1]] : f32 // CHECK-NEXT: spv.Store "Output" [[ARG2]], [[VALUE]] : f32 -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @load_store(%arg0 : !spv.ptr, %arg1 : !spv.ptr) "None" { %1 = spv.Load "Input" %arg0 : f32 spv.Store "Output" %arg1, %1 : f32 @@ -14,7 +14,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @access_chain(%arg0 : !spv.ptr>, Function>, %arg1 : i32, %arg2 : i32) "None" { // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr>, Function> // CHECK-NEXT: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr>, Function> @@ -26,7 +26,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @load_store_zero_rank_float(%arg0: !spv.ptr [0])>, StorageBuffer>, %arg1: !spv.ptr [0])>, StorageBuffer>) "None" { // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr [0])> // CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : f32 @@ -60,7 +60,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @copy_memory_simple() "None" { %0 = spv.Variable : !spv.ptr %1 = spv.Variable : !spv.ptr @@ -72,7 +72,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @copy_memory_different_storage_classes(%in : !spv.ptr, Input>, %out : !spv.ptr, Output>) "None" { // CHECK: spv.CopyMemory "Output" %{{.*}}, "Input" %{{.*}} : !spv.array<4 x f32> spv.CopyMemory "Output" %out, "Input" %in : !spv.array<4xf32> @@ -83,7 +83,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @copy_memory_with_access_operands() "None" { %0 = spv.Variable : !spv.ptr %1 = spv.Variable : !spv.ptr diff --git a/mlir/test/Target/SPIRV/module.mlir b/mlir/test/Target/SPIRV/module.mlir --- a/mlir/test/Target/SPIRV/module.mlir +++ b/mlir/test/Target/SPIRV/module.mlir @@ -1,12 +1,12 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -// CHECK: spv.module Logical GLSL450 requires #spv.vce { +// CHECK: spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK-NEXT: spv.func @foo() "Inline" { // CHECK-NEXT: spv.Return // CHECK-NEXT: } // CHECK-NEXT: } -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "Inline" { spv.Return } @@ -15,18 +15,18 @@ // ----- // CHECK: v1.5 -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { } // ----- // CHECK: [Shader, Float16] -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { } // ----- // CHECK: [SPV_KHR_float_controls, SPV_KHR_subgroup_vote] -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { } diff --git a/mlir/test/Target/SPIRV/non-uniform-ops.mlir b/mlir/test/Target/SPIRV/non-uniform-ops.mlir --- a/mlir/test/Target/SPIRV/non-uniform-ops.mlir +++ b/mlir/test/Target/SPIRV/non-uniform-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @group_non_uniform_ballot spv.func @group_non_uniform_ballot(%predicate: i1) -> vector<4xi32> "None" { // CHECK: %{{.*}} = spv.GroupNonUniformBallot Workgroup %{{.*}}: vector<4xi32> diff --git a/mlir/test/Target/SPIRV/ocl-ops.mlir b/mlir/test/Target/SPIRV/ocl-ops.mlir --- a/mlir/test/Target/SPIRV/ocl-ops.mlir +++ b/mlir/test/Target/SPIRV/ocl-ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Physical64 OpenCL requires #spv.vce { +spv.mlir.module Physical64 OpenCL requires #spv.vce { spv.func @float_insts(%arg0 : f32) "None" { // CHECK: {{%.*}} = spv.OCL.exp {{%.*}} : f32 %0 = spv.OCL.exp %arg0 : f32 diff --git a/mlir/test/Target/SPIRV/phi.mlir b/mlir/test/Target/SPIRV/phi.mlir --- a/mlir/test/Target/SPIRV/phi.mlir +++ b/mlir/test/Target/SPIRV/phi.mlir @@ -2,7 +2,7 @@ // Test branch with one block argument -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: %[[CST:.*]] = spv.Constant 0 %zero = spv.Constant 0 : i32 @@ -23,7 +23,7 @@ // Test branch with multiple block arguments -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: %[[ZERO:.*]] = spv.Constant 0 %zero = spv.Constant 0 : i32 @@ -47,7 +47,7 @@ // Test using block arguments within branch -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: %[[CST0:.*]] = spv.Constant 0 %zero = spv.Constant 0 : i32 @@ -77,7 +77,7 @@ // Test block not following domination order -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: spv.Branch ^bb1 spv.Branch ^bb1 @@ -109,7 +109,7 @@ // Test multiple predecessors -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { %var = spv.Variable : !spv.ptr @@ -158,7 +158,7 @@ // Test nested loops with block arguments -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.GlobalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr, Input> spv.GlobalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr, Input> spv.func @fmul_kernel() "None" { @@ -241,7 +241,7 @@ // Test back-to-back loops with block arguments -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @fmul_kernel() "None" { %cst4 = spv.Constant 4 : i32 diff --git a/mlir/test/Target/SPIRV/sampled-image.mlir b/mlir/test/Target/SPIRV/sampled-image.mlir --- a/mlir/test/Target/SPIRV/sampled-image.mlir +++ b/mlir/test/Target/SPIRV/sampled-image.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: !spv.ptr>, UniformConstant> spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr>, UniformConstant> diff --git a/mlir/test/Target/SPIRV/selection.mlir b/mlir/test/Target/SPIRV/selection.mlir --- a/mlir/test/Target/SPIRV/selection.mlir +++ b/mlir/test/Target/SPIRV/selection.mlir @@ -2,7 +2,7 @@ // Selection with both then and else branches -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @selection(%cond: i1) -> () "None" { // CHECK: spv.Branch ^bb1 // CHECK-NEXT: ^bb1: @@ -55,7 +55,7 @@ // Selection with only then branch // Selection in function entry block -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: spv.func @selection(%[[ARG:.*]]: i1 spv.func @selection(%cond: i1) -> (i32) "None" { // CHECK: spv.Branch ^bb1 diff --git a/mlir/test/Target/SPIRV/spec-constant.mlir b/mlir/test/Target/SPIRV/spec-constant.mlir --- a/mlir/test/Target/SPIRV/spec-constant.mlir +++ b/mlir/test/Target/SPIRV/spec-constant.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: spv.SpecConstant @sc_true = true spv.SpecConstant @sc_true = true // CHECK: spv.SpecConstant @sc_false spec_id(1) = false @@ -48,7 +48,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.SpecConstant @sc_f32_1 = 1.5 : f32 spv.SpecConstant @sc_f32_2 = 2.5 : f32 @@ -68,7 +68,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.SpecConstant @sc_f32_1 = 1.5 : f32 spv.SpecConstant @sc_f32_2 = 2.5 : f32 @@ -88,7 +88,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.SpecConstant @sc_i32_1 = 1 : i32 diff --git a/mlir/test/Target/SPIRV/struct.mlir b/mlir/test/Target/SPIRV/struct.mlir --- a/mlir/test/Target/SPIRV/struct.mlir +++ b/mlir/test/Target/SPIRV/struct.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: !spv.ptr [0])>, Input> spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr [0])>, Input> diff --git a/mlir/test/Target/SPIRV/terminator.mlir b/mlir/test/Target/SPIRV/terminator.mlir --- a/mlir/test/Target/SPIRV/terminator.mlir +++ b/mlir/test/Target/SPIRV/terminator.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK-LABEL: @ret spv.func @ret() -> () "None" { // CHECK: spv.Return diff --git a/mlir/test/Target/SPIRV/undef.mlir b/mlir/test/Target/SPIRV/undef.mlir --- a/mlir/test/Target/SPIRV/undef.mlir +++ b/mlir/test/Target/SPIRV/undef.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { // CHECK: {{%.*}} = spv.undef : f32 // CHECK-NEXT: {{%.*}} = spv.undef : f32 @@ -23,7 +23,7 @@ // ----- -spv.module Logical GLSL450 requires #spv.vce { +spv.mlir.module Logical GLSL450 requires #spv.vce { // CHECK: spv.func {{@.*}} spv.func @ignore_unused_undef() -> () "None" { // CHECK-NEXT: spv.Return diff --git a/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp b/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp --- a/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp +++ b/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp @@ -198,7 +198,8 @@ } 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, diff --git a/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp b/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp --- a/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp +++ b/mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp @@ -44,7 +44,7 @@ }); } - /// 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); }