1//===-- Passes.td - Conversion pass definition file --------*- tablegen -*-===// 2// 3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4// See https://llvm.org/LICENSE.txt for license information. 5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6// 7//===----------------------------------------------------------------------===// 8 9#ifndef MLIR_CONVERSION_PASSES 10#define MLIR_CONVERSION_PASSES 11 12include "mlir/Pass/PassBase.td" 13 14 15//===----------------------------------------------------------------------===// 16// ToLLVM 17//===----------------------------------------------------------------------===// 18 19def ConvertToLLVMPass : Pass<"convert-to-llvm"> { 20 let summary = "Convert to LLVM via dialect interfaces found in the input IR"; 21 let description = [{ 22 This is a generic pass to convert to LLVM, it uses the 23 `ConvertToLLVMPatternInterface` dialect interface to delegate to dialects 24 the injection of conversion patterns. 25 26 If `dynamic` is set to `true`, the pass will look for 27 `ConvertToLLVMAttrInterface` attributes and use them to further configure 28 the conversion process. This option also uses the `DataLayoutAnalysis` 29 analysis to configure the type converter. Enabling this option incurs in 30 extra overhead. 31 }]; 32 33 let constructor = "mlir::createConvertToLLVMPass()"; 34 let options = [ 35 ListOption<"filterDialects", "filter-dialects", "std::string", 36 "Test conversion patterns of only the specified dialects">, 37 Option<"useDynamic", "dynamic", "bool", "false", 38 "Use op conversion attributes to configure the conversion">, 39 ]; 40} 41 42//===----------------------------------------------------------------------===// 43// AffineToStandard 44//===----------------------------------------------------------------------===// 45 46def ConvertAffineToStandard : Pass<"lower-affine"> { 47 let summary = "Lower Affine operations to a combination of Standard and SCF " 48 "operations"; 49 let description = [{ 50 51 Convert operations from the affine dialect into operations from the SCF and 52 standard dialects. 53 54 `affine.for` operations are converted to `scf.for` operations that are free 55 of certain structural restrictions (on their bounds and step). `affine.if` 56 is similarly converted to the `scf.if` operation. `affine.apply` operations 57 are converted into sequences of primitive arithmetic operations from the 58 standard dialect that have the same effect, using operands of the `index` 59 type. Consequently, named maps and sets thare are no longer in use may be 60 removed from the module. 61 62 For example, `%r = affine.apply affine_map<(d0, d1)[s0] -> (d0 + 2*d1 + 63 s0)>(%d0, %d1)[%s0]` 64 can be converted into: 65 66 ```mlir 67 %d0 = <...> 68 %d1 = <...> 69 %s0 = <...> 70 %0 = arith.constant 2 : index 71 %1 = arith.muli %0, %d1 72 %2 = arith.addi %d0, %1 73 %r = arith.addi %2, %s0 74 ``` 75 76 #### Input invariant 77 78 - no `Tensor` types; 79 80 These restrictions may be lifted in the future. 81 82 #### Output IR 83 84 Functions with `affine.for` and `affine.if` operations eliminated. These 85 functions may contain operations from the Standard dialect in addition to 86 those already present before the pass. 87 88 #### Invariants 89 90 - Functions without a body are not modified. 91 - The semantics of the other functions is preserved. 92 - Individual operations other than those mentioned above are not modified 93 if they do not depend on the loop iterator value or on the result of 94 `affine.apply`. 95 }]; 96 let constructor = "mlir::createLowerAffinePass()"; 97 let dependentDialects = [ 98 "memref::MemRefDialect", 99 "scf::SCFDialect", 100 "vector::VectorDialect" 101 ]; 102} 103 104//===----------------------------------------------------------------------===// 105// AMDGPUToROCDL 106//===----------------------------------------------------------------------===// 107 108def ConvertAMDGPUToROCDL : Pass<"convert-amdgpu-to-rocdl"> { 109 let summary = "Convert AMDGPU dialect to ROCDL dialect"; 110 let description = [{ 111 This pass converts supported AMDGPU ops to ROCDL dialect intrinsics. 112 }]; 113 let constructor = "mlir::createConvertAMDGPUToROCDLPass()"; 114 let dependentDialects = [ 115 "LLVM::LLVMDialect", 116 "ROCDL::ROCDLDialect", 117 ]; 118 let options = [Option<"chipset", "chipset", "std::string", 119 /*default=*/"\"gfx000\"", 120 "Chipset that these operations will run on">]; 121} 122 123//===----------------------------------------------------------------------===// 124// ArithToAMDGPU 125//===----------------------------------------------------------------------===// 126def ArithToAMDGPUConversionPass : Pass<"convert-arith-to-amdgpu"> { 127 let summary = "Convert Arith operations to AMDGPU-specific implementations"; 128 let description = [{ 129 Convert `arith` operations (currently extf and truncf on 8-bit floats) 130 to operations in the `amdgpu` dialect. This pass is done in two steps 131 in order to avoid running a notional arith-to-rocdl and arith-to-llvm 132 simultaniously. 133 }]; 134 135 let dependentDialects = ["amdgpu::AMDGPUDialect", "vector::VectorDialect"]; 136 137 let options = [ 138 Option<"chipset", "chipset", "std::string", 139 /*default=*/"\"gfx000\"", 140 "Chipset that these operations will run on">, 141 Option<"saturateFP8Truncf", "saturate-fp8-truncf", "bool", 142 /*default=*/"false", 143 "Use saturating truncation for 8-bit float types">, 144 Option<"allowPackedF16Rtz", "allow-packed-f16-round-to-zero", "bool", 145 /*default=*/"false", 146 "Whether we should allow f32->f16 packed round-to-zero conversion">, 147 ]; 148} 149 150//===----------------------------------------------------------------------===// 151// ArithToEmitC 152//===----------------------------------------------------------------------===// 153 154def ConvertArithToEmitC : Pass<"convert-arith-to-emitc"> { 155 let summary = "Convert Arith dialect to EmitC dialect"; 156 let dependentDialects = ["emitc::EmitCDialect"]; 157} 158 159//===----------------------------------------------------------------------===// 160// ArithToLLVM 161//===----------------------------------------------------------------------===// 162 163def ArithToLLVMConversionPass : Pass<"convert-arith-to-llvm"> { 164 let summary = "Convert Arith dialect to LLVM dialect"; 165 let description = [{ 166 This pass converts supported Arith ops to LLVM dialect instructions. 167 }]; 168 let dependentDialects = ["LLVM::LLVMDialect"]; 169 let options = [ 170 Option<"indexBitwidth", "index-bitwidth", "unsigned", 171 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 172 "Bitwidth of the index type, 0 to use size of machine word">, 173 ]; 174} 175 176//===----------------------------------------------------------------------===// 177// ArithToSPIRV 178//===----------------------------------------------------------------------===// 179 180def ConvertArithToSPIRV : Pass<"convert-arith-to-spirv"> { 181 let summary = "Convert Arith dialect to SPIR-V dialect"; 182 let constructor = "mlir::arith::createConvertArithToSPIRVPass()"; 183 let dependentDialects = ["spirv::SPIRVDialect"]; 184 let options = [ 185 Option<"emulateLT32BitScalarTypes", "emulate-lt-32-bit-scalar-types", 186 "bool", /*default=*/"true", 187 "Emulate narrower scalar types with 32-bit ones if not supported by " 188 "the target">, 189 ]; 190} 191 192//===----------------------------------------------------------------------===// 193// ArithToArmSME 194//===----------------------------------------------------------------------===// 195 196def ArithToArmSMEConversionPass : Pass<"convert-arith-to-arm-sme"> { 197 let summary = "Convert Arith dialect to ArmSME dialect"; 198 let dependentDialects = ["arm_sme::ArmSMEDialect"]; 199} 200 201//===----------------------------------------------------------------------===// 202// ArmNeon2dToIntr 203//===----------------------------------------------------------------------===// 204 205def ConvertArmNeon2dToIntr : Pass<"arm-neon-2d-to-intr"> { 206 let summary = "Convert Arm NEON structured ops to intrinsics"; 207 let constructor = "mlir::createConvertArmNeon2dToIntrPass()"; 208 let dependentDialects = ["arm_neon::ArmNeonDialect", "vector::VectorDialect"]; 209} 210 211//===----------------------------------------------------------------------===// 212// AsyncToLLVM 213//===----------------------------------------------------------------------===// 214 215def ConvertAsyncToLLVMPass : Pass<"convert-async-to-llvm", "ModuleOp"> { 216 let summary = "Convert the operations from the async dialect into the LLVM " 217 "dialect"; 218 let description = [{ 219 Convert `async.execute` operations to LLVM coroutines and use async runtime 220 API to execute them. 221 }]; 222 let dependentDialects = [ 223 "arith::ArithDialect", 224 "async::AsyncDialect", 225 "LLVM::LLVMDialect", 226 "func::FuncDialect", 227 ]; 228} 229 230//===----------------------------------------------------------------------===// 231// BufferizationToMemRef 232//===----------------------------------------------------------------------===// 233 234def ConvertBufferizationToMemRef : Pass<"convert-bufferization-to-memref"> { 235 let summary = "Convert operations from the Bufferization dialect to the " 236 "MemRef dialect"; 237 let description = [{ 238 239 This pass converts bufferization operations into memref operations. 240 241 In the current state, this pass only transforms a `bufferization.clone` 242 operation into `memref.alloc` and `memref.copy` operations and 243 `bufferization.dealloc` operations (the same way as the 244 `-bufferization-lower-deallocations` pass). The conversion of `clone` 245 operations is needed, since some clone operations could remain after 246 applying several transformation processes. Currently, only `canonicalize` 247 transforms clone operations or even eliminates them. This can lead to errors 248 if any clone op survived after all conversion passes (starting from the 249 bufferization dialect) are performed. 250 251 See: 252 https://llvm.discourse.group/t/bufferization-error-related-to-memref-clone/4665 253 254 To avoid these errors, this pass can be performed as a last clean-up pass to 255 transform remaining operations and to proceed in other dialects (memref 256 e.g.). 257 258 Note that this pass only transforms the operation without any further 259 analyses. This pass does not consider any memory analysis or optimization 260 and hence does not resolve any memory leaks. 261 262 }]; 263 let constructor = "mlir::createBufferizationToMemRefPass()"; 264 let dependentDialects = [ 265 "arith::ArithDialect", "memref::MemRefDialect", "scf::SCFDialect", 266 "func::FuncDialect" 267 ]; 268} 269 270//===----------------------------------------------------------------------===// 271// ComplexToLLVM 272//===----------------------------------------------------------------------===// 273 274def ConvertComplexToLLVMPass : Pass<"convert-complex-to-llvm"> { 275 let summary = "Convert Complex dialect to LLVM dialect"; 276 let dependentDialects = ["LLVM::LLVMDialect"]; 277} 278 279//===----------------------------------------------------------------------===// 280// ComplexToLibm 281//===----------------------------------------------------------------------===// 282 283def ConvertComplexToLibm : Pass<"convert-complex-to-libm", "ModuleOp"> { 284 let summary = "Convert Complex dialect to libm calls"; 285 let description = [{ 286 This pass converts supported Complex ops to libm calls. 287 }]; 288 let constructor = "mlir::createConvertComplexToLibmPass()"; 289 let dependentDialects = [ 290 "func::FuncDialect", 291 ]; 292} 293 294//===----------------------------------------------------------------------===// 295// ComplexToSPIRV 296//===----------------------------------------------------------------------===// 297 298def ConvertComplexToSPIRVPass : Pass<"convert-complex-to-spirv"> { 299 let summary = "Convert Complex dialect to SPIRV dialect"; 300 let dependentDialects = ["spirv::SPIRVDialect"]; 301} 302 303//===----------------------------------------------------------------------===// 304// ComplexToStandard 305//===----------------------------------------------------------------------===// 306 307def ConvertComplexToStandard : Pass<"convert-complex-to-standard"> { 308 let summary = "Convert Complex dialect to standard dialect"; 309 let constructor = "mlir::createConvertComplexToStandardPass()"; 310 let dependentDialects = ["math::MathDialect"]; 311} 312 313//===----------------------------------------------------------------------===// 314// ControlFlowToLLVM 315//===----------------------------------------------------------------------===// 316 317def ConvertControlFlowToLLVMPass : Pass<"convert-cf-to-llvm", "ModuleOp"> { 318 let summary = "Convert ControlFlow operations to the LLVM dialect"; 319 let description = [{ 320 Convert ControlFlow operations into LLVM IR dialect operations. 321 322 If other operations are present and their results are required by the LLVM 323 IR dialect operations, the pass will fail. Any LLVM IR operations or types 324 already present in the IR will be kept as is. 325 }]; 326 let dependentDialects = ["LLVM::LLVMDialect"]; 327 let options = [ 328 Option<"indexBitwidth", "index-bitwidth", "unsigned", 329 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 330 "Bitwidth of the index type, 0 to use size of machine word"> 331 ]; 332} 333 334//===----------------------------------------------------------------------===// 335// ControlFlowToSCF 336//===----------------------------------------------------------------------===// 337 338def LiftControlFlowToSCFPass : Pass<"lift-cf-to-scf"> { 339 let summary = "Lift ControlFlow dialect to SCF dialect"; 340 let description = [{ 341 Lifts ControlFlow operations to SCF dialect operations. 342 343 This pass is prefixed with "lift" instead of "convert" as it is not always 344 guaranteed to replace all ControlFlow ops. 345 If a region contains only a single kind of return-like operation, all 346 ControlFlow operations will be replaced successfully. 347 Otherwise a single ControlFlow switch branching to one block per return-like 348 operation kind remains. 349 350 This pass may need to create unreachable terminators in case of infinite 351 loops, which is only supported for 'func.func' for now. If you potentially 352 have infinite loops inside CFG regions not belonging to 'func.func', 353 consider using `transformCFGToSCF` function directly with corresponding 354 `CFGToSCFInterface::createUnreachableTerminator` implementation. 355 }]; 356 357 let dependentDialects = ["scf::SCFDialect", 358 "arith::ArithDialect", 359 "ub::UBDialect", 360 // TODO: This is only necessary until we have a 361 // ub.unreachable op. 362 "func::FuncDialect"]; 363} 364 365//===----------------------------------------------------------------------===// 366// ControlFlowToSPIRV 367//===----------------------------------------------------------------------===// 368 369def ConvertControlFlowToSPIRV : Pass<"convert-cf-to-spirv"> { 370 let summary = "Convert ControlFlow dialect to SPIR-V dialect"; 371 let constructor = "mlir::createConvertControlFlowToSPIRVPass()"; 372 let dependentDialects = ["spirv::SPIRVDialect"]; 373 let options = [ 374 Option<"emulateLT32BitScalarTypes", "emulate-lt-32-bit-scalar-types", 375 "bool", /*default=*/"true", 376 "Emulate narrower scalar types with 32-bit ones if not supported by" 377 " the target"> 378 ]; 379} 380 381//===----------------------------------------------------------------------===// 382// FuncToEmitC 383//===----------------------------------------------------------------------===// 384 385def ConvertFuncToEmitC : Pass<"convert-func-to-emitc", "ModuleOp"> { 386 let summary = "Convert Func dialect to EmitC dialect"; 387 let dependentDialects = ["emitc::EmitCDialect"]; 388} 389 390//===----------------------------------------------------------------------===// 391// FuncToLLVM 392//===----------------------------------------------------------------------===// 393 394def SetLLVMModuleDataLayoutPass : Pass<"set-llvm-module-datalayout", "ModuleOp"> { 395 let summary = "Attach a datalayout string as a module attribute"; 396 let description = [{ 397 Verify that the dataLayout string is a valid LLVM datalayout string and 398 attach it as an attribute `LLVMDialect::getDataLayoutAttrName()` to the 399 module, overriding the existing one. 400 }]; 401 let options = [ 402 Option<"dataLayout", "data-layout", "std::string", 403 /*default=*/"\"\"", 404 "String description (LLVM format) of the data layout that is " 405 "expected on the produced module">, 406 ]; 407} 408 409def ConvertFuncToLLVMPass : Pass<"convert-func-to-llvm", "ModuleOp"> { 410 let summary = "Convert from the Func dialect to the LLVM dialect"; 411 let description = [{ 412 Convert Func dialect operations into the LLVM IR dialect operations. 413 414 #### Input invariant 415 416 - no `tensor` types; 417 - all `vector` are one-dimensional; 418 - all blocks are reachable by following the successors of the first basic 419 block; 420 421 If other operations are present and their results are required by the LLVM 422 IR dialect operations, the pass will fail. Any LLVM IR operations or types 423 already present in the IR will be kept as is. 424 425 An LLVM datalayout string can be attached as an attribute to the module on 426 which the pass anchors. Such an attribute is attached by calling the 427 set-module-datalayout pass. If present, an llvm::DataLayout object is 428 created from this attribute and used in the conversion to LLVM. 429 430 #### Output IR 431 432 Functions converted to LLVM IR. Function arguments types are converted 433 one-to-one. Function results are converted one-to-one and, in case more than 434 1 value is returned, packed into an LLVM IR struct type. Function calls and 435 returns are updated accordingly. Block argument types are updated to use 436 LLVM IR types. 437 }]; 438 let dependentDialects = ["LLVM::LLVMDialect"]; 439 let options = [ 440 Option<"useBarePtrCallConv", "use-bare-ptr-memref-call-conv", "bool", 441 /*default=*/"false", 442 "Replace FuncOp's MemRef arguments with bare pointers to the MemRef " 443 "element types">, 444 Option<"indexBitwidth", "index-bitwidth", "unsigned", 445 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 446 "Bitwidth of the index type, 0 to use size of machine word">, 447 ]; 448} 449 450//===----------------------------------------------------------------------===// 451// FuncToSPIRV 452//===----------------------------------------------------------------------===// 453 454def ConvertFuncToSPIRV : Pass<"convert-func-to-spirv"> { 455 let summary = "Convert Func dialect to SPIR-V dialect"; 456 let constructor = "mlir::createConvertFuncToSPIRVPass()"; 457 let dependentDialects = ["spirv::SPIRVDialect"]; 458 let options = [ 459 Option<"emulateLT32BitScalarTypes", "emulate-lt-32-bit-scalar-types", 460 "bool", /*default=*/"true", 461 "Emulate narrower scalar types with 32-bit ones if not supported by" 462 " the target"> 463 ]; 464} 465 466//===----------------------------------------------------------------------===// 467// GPUCommon 468//===----------------------------------------------------------------------===// 469 470def GpuToLLVMConversionPass : Pass<"gpu-to-llvm", "ModuleOp"> { 471 let summary = "Convert GPU dialect to LLVM dialect with GPU runtime calls"; 472 473 let description = [{ 474 Creates a pass to convert a GPU operations into a sequence of GPU runtime 475 calls. 476 477 This pass does not generate code to call GPU runtime APIs directly but 478 instead uses a small wrapper library that exports a stable and conveniently 479 typed ABI on top of GPU runtimes such as CUDA or ROCm (HIP). 480 }]; 481 482 let options = [ 483 Option<"hostBarePtrCallConv", "use-bare-pointers-for-host", "bool", 484 /*default=*/"false", 485 "Use bare pointers to pass memref arguments to host functions. " 486 "All memrefs must have static shape.">, 487 Option<"kernelBarePtrCallConv", "use-bare-pointers-for-kernels", "bool", 488 /*default=*/"false", 489 "Use bare pointers to pass memref arguments to kernels. " 490 "The kernel must use the same setting for this option." 491 >, 492 Option<"kernelIntersperseSizeCallConv", "intersperse-sizes-for-kernels", "bool", 493 /*default=*/"false", 494 "Inserts a size_t argument following each memref argument, " 495 "containing the static size in bytes of the buffer. Incompatible " 496 "arguments are rejected. This is intended for use by the Vulkan " 497 "runtime with the kernel bare pointer calling convention, to enable " 498 "dynamic binding of buffers as arguments without static type info." 499 > 500 ]; 501 502 let dependentDialects = [ 503 "LLVM::LLVMDialect", 504 "memref::MemRefDialect", 505 ]; 506} 507 508def LowerHostCodeToLLVMPass : Pass<"lower-host-to-llvm", "ModuleOp"> { 509 let summary = "Lowers the host module code and `gpu.launch_func` to LLVM"; 510 511 let description = [{ 512 Creates a pass to emulate `gpu.launch_func` call in LLVM dialect and lower 513 the host module code to LLVM. 514 515 This transformation creates a sequence of global variables that are later 516 linked to the variables in the kernel module, and a series of copies to/from 517 them to emulate the memory transfer from the host or to the device sides. It 518 also converts the remaining Arithmetic, Func, and MemRef dialects into LLVM 519 dialect, emitting C wrappers. 520 }]; 521 522 let dependentDialects = ["LLVM::LLVMDialect"]; 523} 524 525//===----------------------------------------------------------------------===// 526// GPUToLLVMSPV 527//===----------------------------------------------------------------------===// 528 529def ConvertGpuOpsToLLVMSPVOps : Pass<"convert-gpu-to-llvm-spv", "gpu::GPUModuleOp"> { 530 let summary = 531 "Generate LLVM operations to be ingested by a SPIR-V backend for gpu operations"; 532 let dependentDialects = ["LLVM::LLVMDialect"]; 533 let options = [ 534 Option<"use64bitIndex", "use-64bit-index", 535 "bool", /*default=*/"false", 536 "Use 64-bit integers to convert index types">, 537 ]; 538} 539 540//===----------------------------------------------------------------------===// 541// GPUToNVVM 542//===----------------------------------------------------------------------===// 543 544def ConvertGpuOpsToNVVMOps : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> { 545 let summary = "Generate NVVM operations for gpu operations"; 546 let dependentDialects = [ 547 "cf::ControlFlowDialect", 548 "memref::MemRefDialect", 549 "NVVM::NVVMDialect", 550 ]; 551 let options = [ 552 Option<"indexBitwidth", "index-bitwidth", "unsigned", 553 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 554 "Bitwidth of the index type, 0 to use size of machine word">, 555 Option<"hasRedux", "has-redux", "bool", /*default=*/"false", 556 "Target gpu supports redux">, 557 Option<"useBarePtrCallConv", "use-bare-ptr-memref-call-conv", "bool", 558 /*default=*/"false", 559 "Replace memref arguments in GPU functions with bare pointers. " 560 "All memrefs must have static shape."> 561 ]; 562} 563 564//===----------------------------------------------------------------------===// 565// GPUToROCDL 566//===----------------------------------------------------------------------===// 567 568def ConvertGpuOpsToROCDLOps : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> { 569 let summary = "Generate ROCDL operations for gpu operations"; 570 let constructor = "mlir::createLowerGpuOpsToROCDLOpsPass()"; 571 let dependentDialects = [ 572 "ROCDL::ROCDLDialect", 573 "cf::ControlFlowDialect", 574 "memref::MemRefDialect", 575 ]; 576 let options = [ 577 Option<"chipset", "chipset", "std::string", 578 /*default=*/"\"gfx000\"", 579 "Chipset that these operations will run on">, 580 Option<"indexBitwidth", "index-bitwidth", "unsigned", 581 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 582 "Bitwidth of the index type, 0 to use size of machine word">, 583 Option<"useBarePtrCallConv", "use-bare-ptr-memref-call-conv", "bool", 584 /*default=*/"false", 585 "Replace memref arguments in GPU functions with bare pointers." 586 "All memrefs must have static shape">, 587 Option<"runtime", "runtime", "::mlir::gpu::amd::Runtime", 588 "::mlir::gpu::amd::Runtime::Unknown", 589 "Runtime code will be run on (default is Unknown, can also use HIP or OpenCl)", 590 [{::llvm::cl::values( 591 clEnumValN(::mlir::gpu::amd::Runtime::Unknown, "unknown", "Unknown (default)"), 592 clEnumValN(::mlir::gpu::amd::Runtime::HIP, "HIP", "HIP"), 593 clEnumValN(::mlir::gpu::amd::Runtime::OpenCL, "OpenCL", "OpenCL") 594 )}]> 595 ]; 596} 597 598//===----------------------------------------------------------------------===// 599// GPUToSPIRV 600//===----------------------------------------------------------------------===// 601 602def ConvertGPUToSPIRV : Pass<"convert-gpu-to-spirv", "ModuleOp"> { 603 let summary = "Convert GPU dialect to SPIR-V dialect"; 604 let description = [{ 605 This pass converts supported GPU device ops to SPIR-V ops. It does not 606 handle GPU host ops. 607 608 A `gpu.func` op can have parameters to pass in resources. But in SPIR-V 609 entry functions cannot take parameters; they use descriptors to access 610 resources. By default, parameters to a `gpu.func` op will be converted to 611 global variables. These global variables will be assigned sequential binding 612 numbers following their order in the original `gpu.func` op, starting from 613 0, in set 0. One can attach `spirv.interface_var_abi` to those parameters 614 to control the set and binding if wanted. 615 }]; 616 let constructor = "mlir::createConvertGPUToSPIRVPass()"; 617 let dependentDialects = [ 618 "func::FuncDialect", 619 "spirv::SPIRVDialect", 620 ]; 621 let options = [ 622 Option<"use64bitIndex", "use-64bit-index", 623 "bool", /*default=*/"false", 624 "Use 64-bit integers to convert index types">, 625 ]; 626} 627 628//===----------------------------------------------------------------------===// 629// ConvertIndexToLLVMPass 630//===----------------------------------------------------------------------===// 631 632def ConvertIndexToLLVMPass : Pass<"convert-index-to-llvm"> { 633 let summary = "Lower the `index` dialect to the `llvm` dialect."; 634 let description = [{ 635 This pass lowers Index dialect operations to LLVM dialect operations. 636 Operation conversions are 1-to-1 except for the exotic divides: `ceildivs`, 637 `ceildivu`, and `floordivs`, which expand to series of LLVM operations. 638 Importantly, the index bitwidth should be correctly set to the target 639 pointer width via `index-bitwidth`. 640 }]; 641 642 let dependentDialects = ["::mlir::LLVM::LLVMDialect"]; 643 644 let options = [ 645 Option<"indexBitwidth", "index-bitwidth", "unsigned", 646 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 647 "Bitwidth of the index type, 0 to use size of machine word">, 648 ]; 649} 650 651//===----------------------------------------------------------------------===// 652// ConvertIndexToSPIRVPass 653//===----------------------------------------------------------------------===// 654 655def ConvertIndexToSPIRVPass : Pass<"convert-index-to-spirv"> { 656 let summary = "Lower the `index` dialect to the `spirv` dialect."; 657 let description = [{ 658 This pass lowers Index dialect operations to SPIR-V dialect operations. 659 Operation conversions are 1-to-1 except for the exotic divides: `ceildivs`, 660 `ceildivu`, and `floordivs`. The index bitwidth will be 32 or 64 as 661 specified by use-64bit-index. 662 }]; 663 664 let dependentDialects = ["::mlir::spirv::SPIRVDialect"]; 665 666 let options = [ 667 Option<"use64bitIndex", "use-64bit-index", 668 "bool", /*default=*/"false", 669 "Use 64-bit integers to convert index types"> 670 ]; 671} 672 673//===----------------------------------------------------------------------===// 674// LinalgToStandard 675//===----------------------------------------------------------------------===// 676 677def ConvertLinalgToStandard : Pass<"convert-linalg-to-std", "ModuleOp"> { 678 let summary = "Convert the operations from the linalg dialect into the " 679 "Standard dialect"; 680 let constructor = "mlir::createConvertLinalgToStandardPass()"; 681 let dependentDialects = ["func::FuncDialect", "memref::MemRefDialect"]; 682} 683 684//===----------------------------------------------------------------------===// 685// MathToLibm 686//===----------------------------------------------------------------------===// 687 688def ConvertMathToLibm : Pass<"convert-math-to-libm", "ModuleOp"> { 689 let summary = "Convert Math dialect to libm calls"; 690 let description = [{ 691 This pass converts supported Math ops to libm calls. 692 }]; 693 let constructor = "mlir::createConvertMathToLibmPass()"; 694 let dependentDialects = [ 695 "arith::ArithDialect", 696 "func::FuncDialect", 697 "vector::VectorDialect", 698 ]; 699} 700 701//===----------------------------------------------------------------------===// 702// MathToLLVM 703//===----------------------------------------------------------------------===// 704 705def ConvertMathToLLVMPass : Pass<"convert-math-to-llvm"> { 706 let summary = "Convert Math dialect to LLVM dialect"; 707 let dependentDialects = ["LLVM::LLVMDialect"]; 708 let options = [ 709 Option<"approximateLog1p", "approximate-log1p", "bool", "true", 710 "Enable approximation of Log1p."> 711 ]; 712} 713 714//===----------------------------------------------------------------------===// 715// MathToLibm 716//===----------------------------------------------------------------------===// 717 718def ConvertMathToROCDL : Pass<"convert-math-to-rocdl", "ModuleOp"> { 719 let summary = "Convert Math dialect to ROCDL library calls"; 720 let description = [{ 721 This pass converts supported Math ops to ROCDL library calls. 722 }]; 723 let dependentDialects = [ 724 "arith::ArithDialect", 725 "func::FuncDialect", 726 "ROCDL::ROCDLDialect", 727 "vector::VectorDialect", 728 ]; 729} 730 731//===----------------------------------------------------------------------===// 732// MathToSPIRV 733//===----------------------------------------------------------------------===// 734 735def ConvertMathToSPIRV : Pass<"convert-math-to-spirv"> { 736 let summary = "Convert Math dialect to SPIR-V dialect"; 737 let constructor = "mlir::createConvertMathToSPIRVPass()"; 738 let dependentDialects = ["spirv::SPIRVDialect"]; 739} 740 741//===----------------------------------------------------------------------===// 742// MathToEmitC 743//===----------------------------------------------------------------------===// 744 745def ConvertMathToEmitC : Pass<"convert-math-to-emitc"> { 746 let summary = "Convert some Math operations to EmitC call_opaque operations"; 747 let description = [{ 748 This pass converts supported Math ops to `call_opaque` ops targeting libc/libm 749 functions. Unlike convert-math-to-funcs pass, converting to `call_opaque` ops 750 allows to overload the same function with different argument types. 751 }]; 752 let dependentDialects = ["emitc::EmitCDialect"]; 753 let options = [ 754 Option<"languageTarget", "language-target", "::mlir::emitc::LanguageTarget", 755 /*default=*/"::mlir::emitc::LanguageTarget::c99", "Select the language standard target for callees (c99 or cpp11).", 756 [{::llvm::cl::values( 757 clEnumValN(::mlir::emitc::LanguageTarget::c99, "c99", "c99"), 758 clEnumValN(::mlir::emitc::LanguageTarget::cpp11, "cpp11", "cpp11") 759 )}]> 760 ]; 761} 762 763//===----------------------------------------------------------------------===// 764// MathToFuncs 765//===----------------------------------------------------------------------===// 766 767def ConvertMathToFuncs : Pass<"convert-math-to-funcs", "ModuleOp"> { 768 let summary = "Convert Math operations to calls of outlined implementations."; 769 let description = [{ 770 This pass converts supported Math ops to calls of compiler generated 771 functions implementing these operations in software. 772 The LLVM dialect is used for LinkonceODR linkage of the generated functions. 773 }]; 774 let dependentDialects = [ 775 "arith::ArithDialect", 776 "cf::ControlFlowDialect", 777 "func::FuncDialect", 778 "scf::SCFDialect", 779 "vector::VectorDialect", 780 "LLVM::LLVMDialect", 781 ]; 782 let options = [ 783 Option<"minWidthOfFPowIExponent", "min-width-of-fpowi-exponent", "unsigned", 784 /*default=*/"1", 785 "Convert FPowI only if the width of its exponent's integer type " 786 "is greater than or equal to this value">, 787 // Most backend targets support a native ctlz operation, so by default 788 // ctrlz conversion is disabled. 789 Option<"convertCtlz", "convert-ctlz", "bool", /*default=*/"false", 790 "Convert math.ctlz to a software implementation. Enable " 791 "for targets that do not natively support ctlz.">, 792 ]; 793} 794 795//===----------------------------------------------------------------------===// 796// MemRefToEmitC 797//===----------------------------------------------------------------------===// 798 799def ConvertMemRefToEmitC : Pass<"convert-memref-to-emitc"> { 800 let summary = "Convert MemRef dialect to EmitC dialect"; 801 let dependentDialects = ["emitc::EmitCDialect"]; 802} 803 804//===----------------------------------------------------------------------===// 805// MemRefToLLVM 806//===----------------------------------------------------------------------===// 807 808def FinalizeMemRefToLLVMConversionPass : 809 Pass<"finalize-memref-to-llvm", "ModuleOp"> { 810 let summary = "Finalize MemRef dialect to LLVM dialect conversion"; 811 let description = [{ 812 Finalize the conversion of the operations from the MemRef 813 dialect to the LLVM dialect. 814 This conversion will not convert some complex MemRef 815 operations. Make sure to run `expand-strided-metadata` 816 beforehand for these. 817 }]; 818 let dependentDialects = ["LLVM::LLVMDialect"]; 819 let options = [ 820 Option<"useAlignedAlloc", "use-aligned-alloc", "bool", /*default=*/"false", 821 "Use aligned_alloc in place of malloc for heap allocations">, 822 Option<"indexBitwidth", "index-bitwidth", "unsigned", 823 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 824 "Bitwidth of the index type, 0 to use size of machine word">, 825 Option<"useGenericFunctions", "use-generic-functions", 826 "bool", 827 /*default=*/"false", 828 "Use generic allocation and deallocation functions instead of the " 829 "classic 'malloc', 'aligned_alloc' and 'free' functions"> 830 ]; 831} 832 833//===----------------------------------------------------------------------===// 834// MemRefToSPIRV 835//===----------------------------------------------------------------------===// 836 837def MapMemRefStorageClass : Pass<"map-memref-spirv-storage-class"> { 838 let summary = "Map numeric MemRef memory spaces to SPIR-V storage classes"; 839 let constructor = "mlir::createMapMemRefStorageClassPass()"; 840 let dependentDialects = ["spirv::SPIRVDialect"]; 841 let options = [ 842 Option<"clientAPI", "client-api", "std::string", /*default=*/"\"vulkan\"", 843 "The client API to use for populating mappings"> 844 ]; 845} 846 847def ConvertMemRefToSPIRV : Pass<"convert-memref-to-spirv"> { 848 let summary = "Convert MemRef dialect to SPIR-V dialect"; 849 let constructor = "mlir::createConvertMemRefToSPIRVPass()"; 850 let dependentDialects = ["spirv::SPIRVDialect"]; 851 let options = [ 852 Option<"boolNumBits", "bool-num-bits", 853 "int", /*default=*/"8", 854 "The number of bits to store a boolean value">, 855 Option<"use64bitIndex", "use-64bit-index", 856 "bool", /*default=*/"false", 857 "Use 64-bit integers to convert index types"> 858 ]; 859} 860 861//===----------------------------------------------------------------------===// 862// MeshToMPI 863//===----------------------------------------------------------------------===// 864 865def ConvertMeshToMPIPass : Pass<"convert-mesh-to-mpi"> { 866 let summary = "Convert Mesh dialect to MPI dialect."; 867 let constructor = "mlir::createConvertMeshToMPIPass()"; 868 let description = [{ 869 This pass converts communication operations from the Mesh dialect to the 870 MPI dialect. 871 If it finds a global named "static_mpi_rank" it will use that splat value 872 instead of calling MPI_Comm_rank. This allows optimizations like constant 873 shape propagation and fusion because shard/partition sizes depend on the 874 rank. 875 }]; 876 let dependentDialects = [ 877 "memref::MemRefDialect", 878 "mpi::MPIDialect", 879 "scf::SCFDialect", 880 "bufferization::BufferizationDialect" 881 ]; 882} 883 884//===----------------------------------------------------------------------===// 885// NVVMToLLVM 886//===----------------------------------------------------------------------===// 887 888def ConvertNVVMToLLVMPass : Pass<"convert-nvvm-to-llvm"> { 889 let summary = "Convert NVVM to PTX with Inline Assembly in LLVM dialect"; 890 let description = [{ 891 This pass generates PTX instructions using inline assembly for NVVM 892 operations implements `BasicPtxBuilderInterface`. 893 }]; 894 let dependentDialects = [ 895 "NVVM::NVVMDialect", 896 ]; 897} 898 899//===----------------------------------------------------------------------===// 900// NVGPUToNVVM 901//===----------------------------------------------------------------------===// 902 903def ConvertNVGPUToNVVMPass : Pass<"convert-nvgpu-to-nvvm"> { 904 let summary = "Convert NVGPU dialect to NVVM dialect"; 905 let description = [{ 906 This pass converts supported NVGPU ops to NVVM dialect intrinsics. 907 }]; 908 909 let dependentDialects = [ 910 "NVVM::NVVMDialect", 911 ]; 912} 913 914 915//===----------------------------------------------------------------------===// 916// OpenACCToSCF 917//===----------------------------------------------------------------------===// 918 919def ConvertOpenACCToSCF : Pass<"convert-openacc-to-scf", "ModuleOp"> { 920 let summary = "Convert the OpenACC ops to OpenACC with SCF dialect"; 921 let constructor = "mlir::createConvertOpenACCToSCFPass()"; 922 let dependentDialects = ["scf::SCFDialect", "acc::OpenACCDialect"]; 923} 924 925//===----------------------------------------------------------------------===// 926// OpenMPToLLVM 927//===----------------------------------------------------------------------===// 928 929def ConvertOpenMPToLLVMPass : Pass<"convert-openmp-to-llvm", "ModuleOp"> { 930 let summary = "Convert the OpenMP ops to OpenMP ops with LLVM dialect"; 931 let dependentDialects = ["LLVM::LLVMDialect"]; 932} 933 934//===----------------------------------------------------------------------===// 935// PDLToPDLInterp 936//===----------------------------------------------------------------------===// 937 938def ConvertPDLToPDLInterp : Pass<"convert-pdl-to-pdl-interp", "ModuleOp"> { 939 let summary = "Convert PDL ops to PDL interpreter ops"; 940 let constructor = "mlir::createPDLToPDLInterpPass()"; 941 let dependentDialects = ["pdl_interp::PDLInterpDialect"]; 942} 943 944//===----------------------------------------------------------------------===// 945// ReconcileUnrealizedCasts 946//===----------------------------------------------------------------------===// 947 948def ReconcileUnrealizedCasts : Pass<"reconcile-unrealized-casts"> { 949 let summary = "Simplify and eliminate unrealized conversion casts"; 950 let description = [{ 951 Eliminate `unrealized_conversion_cast` operations, commonly introduced by 952 partial dialect conversions, that transitively convert a value to another 953 value of the same type, that is: 954 955 ``` 956 %0 = "producer.op"() : () -> !type.A 957 %1 = unrealized_conversion_cast %0 : !type.A to !type.B 958 %2 = unrealized_conversion_cast %1 : !type.B to !type.C 959 %3 = unrealized_conversion_cast %2 : !type.C to !type.A 960 "consumer.op"(%3) : (!type.A) -> () 961 ``` 962 963 Such situations appear when the consumer operation is converted by one pass 964 and the producer operation is converted by another pass, each of which 965 produces an unrealized cast. This pass can be used to clean up the IR. 966 }]; 967 let constructor = "mlir::createReconcileUnrealizedCastsPass()"; 968} 969 970//===----------------------------------------------------------------------===// 971// SCFToControlFlow 972//===----------------------------------------------------------------------===// 973 974def SCFToControlFlow : Pass<"convert-scf-to-cf"> { 975 let summary = "Convert SCF dialect to ControlFlow dialect, replacing structured" 976 " control flow with a CFG"; 977 let constructor = "mlir::createConvertSCFToCFPass()"; 978 let dependentDialects = ["cf::ControlFlowDialect"]; 979} 980 981//===----------------------------------------------------------------------===// 982// SCFToOpenMP 983//===----------------------------------------------------------------------===// 984 985def ConvertSCFToOpenMPPass : Pass<"convert-scf-to-openmp", "ModuleOp"> { 986 let summary = "Convert SCF parallel loop to OpenMP parallel + workshare " 987 "constructs."; 988 989 let options = [ 990 Option<"numThreads", "num-threads", "unsigned", 991 /*default=kUseOpenMPDefaultNumThreads*/"0", 992 "Number of threads to use"> 993 ]; 994 995 let dependentDialects = ["omp::OpenMPDialect", "LLVM::LLVMDialect", 996 "memref::MemRefDialect"]; 997} 998 999//===----------------------------------------------------------------------===// 1000// SCFToSPIRV 1001//===----------------------------------------------------------------------===// 1002 1003def SCFToSPIRV : Pass<"convert-scf-to-spirv"> { 1004 let summary = "Convert SCF dialect to SPIR-V dialect."; 1005 let description = [{ 1006 Converts SCF ops into SPIR-V structured control flow ops. 1007 SPIR-V structured control flow ops do not support yielding values. 1008 So for SCF ops yielding values, SPIR-V variables are created for 1009 holding the values and load/store operations are emitted for updating 1010 them. 1011 }]; 1012 let dependentDialects = ["spirv::SPIRVDialect"]; 1013} 1014 1015//===----------------------------------------------------------------------===// 1016// SCFToGPU 1017//===----------------------------------------------------------------------===// 1018 1019def ConvertAffineForToGPU 1020 : InterfacePass<"convert-affine-for-to-gpu", "FunctionOpInterface"> { 1021 let summary = "Convert top-level AffineFor Ops to GPU kernels"; 1022 let constructor = "mlir::createAffineForToGPUPass()"; 1023 let dependentDialects = ["gpu::GPUDialect"]; 1024 let options = [ 1025 Option<"numBlockDims", "gpu-block-dims", "unsigned", /*default=*/"1u", 1026 "Number of GPU block dimensions for mapping">, 1027 Option<"numThreadDims", "gpu-thread-dims", "unsigned", /*default=*/"1u", 1028 "Number of GPU thread dimensions for mapping"> 1029 ]; 1030} 1031 1032def ConvertParallelLoopToGpu : Pass<"convert-parallel-loops-to-gpu"> { 1033 let summary = "Convert mapped scf.parallel ops to gpu launch operations"; 1034 let constructor = "mlir::createParallelLoopToGpuPass()"; 1035 let dependentDialects = ["affine::AffineDialect", "gpu::GPUDialect"]; 1036} 1037 1038//===----------------------------------------------------------------------===// 1039// SCFToEmitC 1040//===----------------------------------------------------------------------===// 1041 1042def SCFToEmitC : Pass<"convert-scf-to-emitc"> { 1043 let summary = "Convert SCF dialect to EmitC dialect, maintaining structured" 1044 " control flow"; 1045 let dependentDialects = ["emitc::EmitCDialect"]; 1046} 1047 1048//===----------------------------------------------------------------------===// 1049// ShapeToStandard 1050//===----------------------------------------------------------------------===// 1051 1052def ConvertShapeToStandard : Pass<"convert-shape-to-std", "ModuleOp"> { 1053 let summary = "Convert operations from the shape dialect into the standard " 1054 "dialect"; 1055 let constructor = "mlir::createConvertShapeToStandardPass()"; 1056 let dependentDialects = [ 1057 "scf::SCFDialect", 1058 ]; 1059} 1060 1061def ConvertShapeConstraints : Pass<"convert-shape-constraints"> { 1062 let summary = "Convert shape constraint operations to the standard dialect"; 1063 let description = [{ 1064 This pass eliminates shape constraints from the program, converting them to 1065 eager (side-effecting) error handling code. 1066 1067 This pass is separate from the regular convert-shape-to-standard, despite 1068 converting between the same dialects, because converting shape constraints 1069 can happen at a different part of the program than general shape 1070 computation lowering. 1071 }]; 1072 let constructor = "mlir::createConvertShapeConstraintsPass()"; 1073 let dependentDialects = ["cf::ControlFlowDialect", "scf::SCFDialect"]; 1074} 1075 1076//===----------------------------------------------------------------------===// 1077// SPIRVToLLVM 1078//===----------------------------------------------------------------------===// 1079 1080def ConvertSPIRVToLLVMPass : Pass<"convert-spirv-to-llvm", "ModuleOp"> { 1081 let summary = "Convert SPIR-V dialect to LLVM dialect"; 1082 let description = [{ 1083 See https://mlir.llvm.org/docs/SPIRVToLLVMDialectConversion/ 1084 for more details. 1085 }]; 1086 let dependentDialects = ["LLVM::LLVMDialect"]; 1087 1088 let options = [ 1089 Option<"clientAPI", "client-api", "::mlir::spirv::ClientAPI", 1090 /*default=*/"::mlir::spirv::ClientAPI::Unknown", 1091 "Derive StorageClass to address space mapping from the client API", 1092 [{::llvm::cl::values( 1093 clEnumValN(::mlir::spirv::ClientAPI::Unknown, "Unknown", "Unknown (default)"), 1094 clEnumValN(::mlir::spirv::ClientAPI::Metal, "Metal", "Metal"), 1095 clEnumValN(::mlir::spirv::ClientAPI::OpenCL, "OpenCL", "OpenCL"), 1096 clEnumValN(::mlir::spirv::ClientAPI::Vulkan, "Vulkan", "Vulkan"), 1097 clEnumValN(::mlir::spirv::ClientAPI::WebGPU, "WebGPU", "WebGPU") 1098 )}]>, 1099 ]; 1100} 1101 1102//===----------------------------------------------------------------------===// 1103// TensorToLinalg 1104//===----------------------------------------------------------------------===// 1105 1106def ConvertTensorToLinalg : Pass<"convert-tensor-to-linalg", "ModuleOp"> { 1107 let summary = "Convert some Tensor dialect ops to Linalg dialect"; 1108 let constructor = "mlir::createConvertTensorToLinalgPass()"; 1109 let dependentDialects = [ 1110 "arith::ArithDialect", 1111 "linalg::LinalgDialect", 1112 ]; 1113} 1114 1115//===----------------------------------------------------------------------===// 1116// TensorToSPIRV 1117//===----------------------------------------------------------------------===// 1118 1119def ConvertTensorToSPIRV : Pass<"convert-tensor-to-spirv"> { 1120 let summary = "Convert Tensor dialect to SPIR-V dialect"; 1121 let constructor = "mlir::createConvertTensorToSPIRVPass()"; 1122 let dependentDialects = ["spirv::SPIRVDialect"]; 1123 let options = [ 1124 Option<"emulateLT32BitScalarTypes", "emulate-lt-32-bit-scalar-types", 1125 "bool", /*default=*/"true", 1126 "Emulate narrower scalar types with 32-bit ones if not supported by" 1127 " the target"> 1128 ]; 1129} 1130 1131//===----------------------------------------------------------------------===// 1132// TosaToArith 1133//===----------------------------------------------------------------------===// 1134 1135def TosaToArith : Pass<"tosa-to-arith"> { 1136 let summary = "Lower TOSA to the Arith dialect"; 1137 let dependentDialects = [ 1138 "arith::ArithDialect", 1139 ]; 1140 let description = [{ 1141 Pass that converts TOSA operations to the equivalent operations using the 1142 operations in the Arith dialect. The ApplyScale operator is optionally 1143 included as it is often preserved until the final invocation. 1144 }]; 1145 1146 let options = [ 1147 Option<"includeApplyRescale", "include-apply-rescale", 1148 "bool", /*default=*/"false", 1149 "Whether to include the lowering for tosa.apply_rescale to arith">, 1150 Option<"use32Bit", "use-32-bit", 1151 "bool", /*default=*/"false", 1152 "Whether to prioritze lowering to 32-bit operations"> 1153 ]; 1154 1155 let constructor = "tosa::createTosaToArith()"; 1156} 1157 1158//===----------------------------------------------------------------------===// 1159// TosaToLinalg 1160//===----------------------------------------------------------------------===// 1161 1162def TosaToLinalg 1163 : InterfacePass<"tosa-to-linalg", "FunctionOpInterface"> { 1164 let summary = "Lower TOSA to LinAlg on tensors"; 1165 let description = [{ 1166 Pass that converts TOSA operations to the equivalent operations using the 1167 tensor operations in LinAlg. 1168 }]; 1169 1170 let constructor = "tosa::createTosaToLinalg()"; 1171 let options = [ 1172 Option<"disableTosaDecompositions", "disable-tosa-decompositions", 1173 "bool", /*default=*/"false", 1174 "Disable tosa decompositions pass">, 1175 Option<"aggressiveReduceConstant", "aggressive-reduce-constant", 1176 "bool", /*default=*/"false", 1177 "Always perform the reduce constant optimization"> 1178 ]; 1179} 1180 1181//===----------------------------------------------------------------------===// 1182// TosaToLinalgNamed 1183//===----------------------------------------------------------------------===// 1184 1185def TosaToLinalgNamed 1186 : InterfacePass<"tosa-to-linalg-named", "FunctionOpInterface"> { 1187 let summary = "Lower TOSA to LinAlg named operations"; 1188 let description = [{ 1189 Pass that converts TOSA operations to the equivalent operations using the 1190 Linalg named operations. 1191 }]; 1192 1193 let options = [ 1194 Option<"preferConv2DKernelLayoutHWCF", "prefer-conv2d-kernel-layout-hwcf", 1195 "bool", /*default=*/"false", 1196 "Prefer generating linalg.conv_2d_nhwc_hwcf over linalg.conv_2d_nhwc_fhwc"> 1197 ]; 1198 1199 let constructor = "tosa::createTosaToLinalgNamed()"; 1200} 1201 1202//===----------------------------------------------------------------------===// 1203// TosaToMLProgram 1204//===----------------------------------------------------------------------===// 1205 1206def TosaToMLProgram : Pass<"tosa-to-mlprogram", "ModuleOp"> { 1207 let summary = "Lower TOSA to the MLProgram dialect"; 1208 let dependentDialects = ["ml_program::MLProgramDialect"]; 1209 let description = [{ 1210 Pass that converts TOSA's variable operator operations to the equivalent 1211 MLProgram operations. 1212 }]; 1213} 1214 1215//===----------------------------------------------------------------------===// 1216// TosaToSCF 1217//===----------------------------------------------------------------------===// 1218 1219def TosaToSCF : Pass<"tosa-to-scf"> { 1220 let summary = "Lower TOSA to the SCF dialect"; 1221 let dependentDialects = ["tensor::TensorDialect, scf::SCFDialect"]; 1222 let description = [{ 1223 Pass that converts TOSA's control flow operations to the equivalent SCF 1224 operations. 1225 }]; 1226 1227 let constructor = "tosa::createTosaToSCF()"; 1228} 1229 1230//===----------------------------------------------------------------------===// 1231// TosaToTensor 1232//===----------------------------------------------------------------------===// 1233 1234def TosaToTensor : Pass<"tosa-to-tensor"> { 1235 let summary = "Lower TOSA to the Tensor dialect"; 1236 let dependentDialects = [ 1237 "tensor::TensorDialect", 1238 ]; 1239 let description = [{ 1240 Pass that converts TOSA operations to the equivalent operations using the 1241 operations in the Tensor dialect. 1242 }]; 1243 1244 let constructor = "tosa::createTosaToTensor()"; 1245} 1246 1247//===----------------------------------------------------------------------===// 1248// UBToLLVM 1249//===----------------------------------------------------------------------===// 1250 1251def UBToLLVMConversionPass : Pass<"convert-ub-to-llvm"> { 1252 let summary = "Convert UB dialect to LLVM dialect"; 1253 let description = [{ 1254 This pass converts supported UB ops to LLVM dialect instructions. 1255 }]; 1256 let dependentDialects = ["LLVM::LLVMDialect"]; 1257 let options = [ 1258 Option<"indexBitwidth", "index-bitwidth", "unsigned", 1259 /*default=kDeriveIndexBitwidthFromDataLayout*/"0", 1260 "Bitwidth of the index type, 0 to use size of machine word">, 1261 ]; 1262} 1263 1264//===----------------------------------------------------------------------===// 1265// UBToSPIRV 1266//===----------------------------------------------------------------------===// 1267 1268def UBToSPIRVConversionPass : Pass<"convert-ub-to-spirv"> { 1269 let summary = "Convert UB dialect to SPIR-V dialect"; 1270 let description = [{ 1271 This pass converts supported UB ops to SPIR-V dialect ops. 1272 }]; 1273 let dependentDialects = ["spirv::SPIRVDialect"]; 1274} 1275 1276//===----------------------------------------------------------------------===// 1277// VectorToGPU 1278//===----------------------------------------------------------------------===// 1279 1280def ConvertVectorToGPU : Pass<"convert-vector-to-gpu"> { 1281 let summary = "Lower the operations from the vector dialect into the GPU " 1282 "dialect"; 1283 let constructor = "mlir::createConvertVectorToGPUPass()"; 1284 let dependentDialects = [ 1285 "memref::MemRefDialect", "gpu::GPUDialect", "affine::AffineDialect", 1286 "vector::VectorDialect", "nvgpu::NVGPUDialect" 1287 ]; 1288 1289 let options = [ 1290 Option<"useNvGpu", "use-nvgpu", "bool", /*default=*/"false", 1291 "convert to NvGPU ops instead of GPU dialect ops"> 1292 ]; 1293} 1294 1295//===----------------------------------------------------------------------===// 1296// VectorToSCF 1297//===----------------------------------------------------------------------===// 1298 1299def ConvertVectorToSCF : Pass<"convert-vector-to-scf"> { 1300 let summary = "Lower the operations from the vector dialect into the SCF " 1301 "dialect"; 1302 let constructor = "mlir::createConvertVectorToSCFPass()"; 1303 let dependentDialects = [ 1304 "affine::AffineDialect", 1305 "memref::MemRefDialect", 1306 "scf::SCFDialect", 1307 "tensor::TensorDialect" 1308 ]; 1309 let options = [ 1310 Option<"fullUnroll", "full-unroll", "bool", /*default=*/"false", 1311 "Perform full unrolling when converting vector transfers to SCF">, 1312 Option<"targetRank", "target-rank", "unsigned", /*default=*/"1", 1313 "Target vector rank to which transfer ops should be lowered">, 1314 Option<"lowerTensors", "lower-tensors", "bool", /*default=*/"false", 1315 "Lower transfer ops that operate on tensors">, 1316 Option<"lowerScalable", "lower-scalable", "bool", /*default=*/"false", 1317 "Add scalable vector specific lowerings (that introduce loops)"> 1318 ]; 1319} 1320 1321//===----------------------------------------------------------------------===// 1322// VectorToArmSME 1323//===----------------------------------------------------------------------===// 1324 1325def ConvertVectorToArmSME : Pass<"convert-vector-to-arm-sme"> { 1326 let summary = "Lower the operations from the vector dialect into the ArmSME " 1327 "dialect"; 1328 let constructor = "mlir::createConvertVectorToArmSMEPass()"; 1329 let description = [{ 1330 Pass that converts vector dialect operations into equivalent ArmSME dialect 1331 operations. 1332 }]; 1333 let dependentDialects = ["arm_sme::ArmSMEDialect", "arm_sve::ArmSVEDialect"]; 1334} 1335 1336//===----------------------------------------------------------------------===// 1337// ArmSMEToSCF 1338//===----------------------------------------------------------------------===// 1339 1340def ConvertArmSMEToSCF : Pass<"convert-arm-sme-to-scf"> { 1341 let summary = "Lower the operations from the ArmSME dialect into the SCF " 1342 "dialect"; 1343 let constructor = "mlir::createConvertArmSMEToSCFPass()"; 1344 let dependentDialects = [ 1345 "scf::SCFDialect", 1346 "arith::ArithDialect", 1347 "vector::VectorDialect", 1348 "arm_sme::ArmSMEDialect" 1349 ]; 1350} 1351 1352//===----------------------------------------------------------------------===// 1353// ArmSMEToLLVM 1354//===----------------------------------------------------------------------===// 1355 1356def ConvertArmSMEToLLVM : InterfacePass<"convert-arm-sme-to-llvm", "FunctionOpInterface"> { 1357 let summary = "Lower the operations from the ArmSME dialect into the LLVM " 1358 "dialect"; 1359 let constructor = "mlir::createConvertArmSMEToLLVMPass()"; 1360 let dependentDialects = [ 1361 "arm_sme::ArmSMEDialect", 1362 "LLVM::LLVMDialect" 1363 ]; 1364 let options = [ 1365 Option<"dumpTileLiveRanges", "dump-tile-live-ranges", 1366 "bool", /*default=*/"false", 1367 "Dump the live ranges of SME tiles (for debugging)"> 1368 ]; 1369} 1370 1371//===----------------------------------------------------------------------===// 1372// VectorToLLVM 1373//===----------------------------------------------------------------------===// 1374 1375def ConvertVectorToLLVMPass : Pass<"convert-vector-to-llvm"> { 1376 let summary = "Lower the operations from the vector dialect into the LLVM " 1377 "dialect"; 1378 let description = [{ 1379 1380 Convert operations from the vector dialect into the LLVM IR dialect 1381 operations. The lowering pass provides several options to control 1382 the kinds of optimizations that are allowed. It also provides options 1383 that enable the use of one or more architectural-specific dialects 1384 (AMX, X86Vector, ArmNeon, ArmSVE, etc.) in combination with the 1385 architectural-neutral vector dialect lowering. 1386 1387 }]; 1388 // Override explicitly in C++ to allow conditional dialect dependence. 1389 // let dependentDialects; 1390 let options = [ 1391 Option<"reassociateFPReductions", "reassociate-fp-reductions", 1392 "bool", /*default=*/"false", 1393 "Allows llvm to reassociate floating-point reductions for speed">, 1394 Option<"force32BitVectorIndices", "force-32bit-vector-indices", 1395 "bool", /*default=*/"true", 1396 "Allows compiler to assume vector indices fit in 32-bit if that " 1397 "yields faster code">, 1398 Option<"amx", "enable-amx", 1399 "bool", /*default=*/"false", 1400 "Enables the use of AMX dialect while lowering the vector " 1401 "dialect.">, 1402 Option<"armNeon", "enable-arm-neon", 1403 "bool", /*default=*/"false", 1404 "Enables the use of ArmNeon dialect while lowering the vector " 1405 "dialect.">, 1406 Option<"armSVE", "enable-arm-sve", 1407 "bool", /*default=*/"false", 1408 "Enables the use of ArmSVE dialect while lowering the vector " 1409 "dialect.">, 1410 Option<"x86Vector", "enable-x86vector", 1411 "bool", /*default=*/"false", 1412 "Enables the use of X86Vector dialect while lowering the vector " 1413 "dialect.">, 1414 Option<"vectorTransformsOptions", "vector-transform-options", 1415 "vector::VectorTransformsOptions", 1416 /*default=*/"vector::VectorTransformsOptions()", 1417 "Options to lower some operations like contractions and transposes.">, 1418 ]; 1419} 1420 1421//===----------------------------------------------------------------------===// 1422// VectorToSPIRV 1423//===----------------------------------------------------------------------===// 1424 1425def ConvertVectorToSPIRV : Pass<"convert-vector-to-spirv"> { 1426 let summary = "Convert Vector dialect to SPIR-V dialect"; 1427 let constructor = "mlir::createConvertVectorToSPIRVPass()"; 1428 let dependentDialects = [ 1429 "spirv::SPIRVDialect", 1430 "ub::UBDialect" 1431 ]; 1432} 1433 1434//===----------------------------------------------------------------------===// 1435// VectorToXeGPU 1436//===----------------------------------------------------------------------===// 1437 1438def ConvertVectorToXeGPU : Pass<"convert-vector-to-xegpu"> { 1439 let summary = "Lower the operations from the vector dialect into the XeGPU " 1440 "dialect"; 1441 let constructor = "mlir::createConvertVectorToXeGPUPass()"; 1442 let dependentDialects = [ 1443 "memref::MemRefDialect", "arith::ArithDialect", 1444 "vector::VectorDialect", "xegpu::XeGPUDialect" 1445 ]; 1446} 1447 1448#endif // MLIR_CONVERSION_PASSES 1449